高性能计算通信深度学习分布式训练【免费下载链接】ncclOptimized primitives for collective multi-GPU communication项目地址https://gitcode.com/gh_mirrors/nc/nccl点击查看免费下载GINGPU-Initiated Networking是 NCCL Device API 中用于跨节点网络通信的核心模块它允许设备GPU内核直接发起网络数据传输而无需 CPU 介入。本文以仓库中 GIN 设备 API 参考文档 为主体结合 Device API 使用指南 与仓库源码gin.h、gin_device_common.h、gin_barrier.h以及 纯 GIN AlltoAll 示例系统讲解ncclGin对象、单向put/get传输、信号Signal与计数器Counter完成机制、以及ncclGinBarrierSession网络屏障的完整用法。读完本文你将能够在自己的 CUDA 内核中编写纯 GIN 的单向通信与跨节点同步逻辑。GIN 是什么从 Device API 到 GPU 发起网络NCCL 自 2.28 版本起提供设备端通信 APIDevice API使通信原语可以直接在用户 CUDA 内核中使用而 GIN 自NCCL 2.28.7起成为其中负责跨网络通信的模块见 deviceapi.rst。与 LSALoad/Store Accessible面向 NVLink/PCIe P2P 的直访通信和 MultimemNVLink SHARP 硬件多播不同GIN 的所有数据传输都通过网卡完成适用于不同节点之间、无法使用 LSA 直访的 rank 之间的通信纯网络路径的集体通信如跨节点 AlltoAll需要计算与通信重叠、以信号驱动异步完成检测的场景。从源码结构看GIN 在设备侧由 gin.h 中的ncclGin类型与 gin_device_common.h 中的后端分发机制支撑支持三种后端见NCCL_GIN_BACKEND_MASK_ALL宏与 device_setup.rst 中ncclGinType_t的说明Host ProxyNCCL_GIN_TYPE_PROXY由主机代理转发的 GIN 类型GDAKINCCL_GIN_TYPE_GDAKIGPUDirect Async Kernel-Initiated内核直接发起GPINCCL_GIN_TYPE_GPIGPU-Push Interface需要 SpectrumX自 NCCL 2.30.6 起作为实验特性加入。使用 GIN 的软硬件要求根据 deviceapi.rst 的Requirements一节使用 GIN 需要满足编译 GPU 代码需 CUDA 12.2 或更高GPU 需 Volta 及以上驱动 510.40.3网卡需 CX4 及以上rdma-core 44.0主机代理 GIN 需要 DMA-BUF 或 nvidia-peermemGDAKI 需要 DMA-BUF内核 6.1或 nvidia-peermem网络拓扑要求全 NIC 连通不支持跨 rail 无法通信的拓扑也不支持NCCL_CROSS_NIC0不支持融合网卡Fused NIC双口网卡需设置NCCL_IB_MERGE_NICS0使用由多个 cuMem 段支撑的缓冲区时需要 DMA-BUF。此外跨版本兼容性方面LSA/Multimem 内核自 NCCL 2.29 起支持向后兼容而GIN 内核目前不做向后兼容NCCL 升级后需要重新编译。ncclGin一个对象对应一条 GIN 上下文ncclGin是 GIN 支持的主类封装了 GIN 的关键元素文档中定义为A class encompassing major elements of the GIN support。它只在设备GPU代码中可调用构造方式如下ncclGin(ncclDevComm const comm, int contextIndex)comm通过ncclDevCommCreate创建的设备通信器参见 device_setup.rstcontextIndexGIN 上下文网络通信通道的索引。多个 GIN 上下文可以将流量分散到多条连接上避免锁竞争和瓶颈因此面向性能的内核应当在不同上下文间轮转使用以提高资源利用率。可用上下文数量通过设备通信器的ginContextCount字段查询该字段在 device_setup.rst 中注明自 NCCL 2.28.7 起提供。在 core.h 中可以看到ncclGinSignal_t与ncclGinCounter_t均为 32 位无符号整数句柄而ncclGinRequest_t是一个 16 字节两个uint64_t的不透明请求对象用于异步flushAsync/wait。ncclGin的一个实例总是代表某一个特定上下文。当需要执行一个栅栏barrier且该栅栏的 fence 需要覆盖通信器上的所有GIN 上下文时例如操作被切分到多个上下文、多网卡场景可以在ncclGinBarrier的位置上传入ncclGinAllContexts(comm)而非某个ncclGin。ncclGinAllContexts在 gin_barrier.h 中定义其注释明确说明传给它时信号仍走上下文 0而 fence 会遍历通信器上的每一个上下文。从 gin.h 的实现看ncclGin实际上是ncclGin_BackendMask...的别名using ncclGin ncclGin_BackendMaskNCCL_GIN_BACKEND_MASK_ALL其中NCCL_GIN_BACKEND_MASK_ALL按编译期宏NCCL_GIN_PROXY_ENABLE、NCCL_GIN_GDAKI_ENABLE、NCCL_GIN_GPI_ENABLE展开为后端位掩码见 gin_device_common.h。运行时通过ncclGinCall/ncclGinCallImpl按位掩码分发到对应后端的ncclGinApi_Put、ncclGinApi_Get、ncclGinApi_Flush等实现——这正是一个 API、多后端的设计。单向数据传输put 与 getput本地缓冲区 → 远端缓冲区void put(ncclTeam team, int peer, ncclWindow_t dstWnd, size_t dstOffset, ncclWindow_t srcWnd, size_t srcOffset, size_t bytes, RemoteAction remoteAction, LocalAction localAction, Coop coop, DescriptorSmem descriptor, cuda::thread_scope alreadyReleased, cuda::thread_scope expected_scope, SegmentType bufType)put调度一个由设备发起的单边one-sided数据传输操作把本地缓冲区的内容写到 peer 的远端缓冲区。参数要点peer是team内的一个 rank参见设备 API 的 Teams 概念ncclTeamWorld()、ncclTeamLsa()、ncclTeamRail()它可以指向本地 rank 自身即 loopback 回环目的与源缓冲区各自由窗口dstWnd/srcWnd 字节偏移dstOffset/srcOffset描述bytes指定传输的字节数若 GIN 以NCCL_GIN_CONNECTION_RAIL连接类型初始化peer必须与本地 rank 处于同一个 rail team 内rail team 是本地 LSA team 内 rank 编号相同的对等节点集合见 deviceapi.rst 的 Teams 一节。前七个参数之外均为可选参数参数作用默认值remoteAction载荷在目的端落定后在目的 peer 上执行的动作ncclGin_None无动作localAction输入被消费后在本地 rank 上执行的动作ncclGin_Nonecoop参与本次操作的线程集合ncclCoopThread单个设备线程官方推荐模型descriptor可选的共享内存描述符空间ncclGin_NonealreadyReleased/expected_scope线程作用域thread scope控制thread_scope_thread/thread_scope_devicebufType源/目的缓冲区的物理内存组成ncclGin_SegmentDeviceremoteAction的常见取值是ncclGin_Signal{Inc|Add}信号类动作localAction的取值是ncclGin_CounterInc计数器动作详见下文信号与计数器一节。bufType的取值见 deviceapi.rst 的 Segment Types 一节ncclGin_SegmentDevice默认纯设备显存段、ncclGin_SegmentHostNuma纯 CPU 侧HOST_NUMA段、ncclGin_SegmentMixed设备段与HOST_NUMA段混合。可见性语义目的端 peer 上信号的可见性意味着它所附着的那次put数据也已可见。根据信号类型的不同强/弱见下文信号的可见性还可能意味着同一上下文上、发往同一 peer 的所有此前 put 数据均已可见。ncclGin还定义了两种便捷变体ncclSymPtr变体用ncclSymPtrT类型描述缓冲区且大小以元素个数nElts而非字节数表达putValue变体一次传输单个元素不超过 8 字节按值传递无需窗口偏移寻址。get远端缓冲区 → 本地缓冲区void get(ncclTeam team, int peer, ncclWindow_t remoteWnd, size_t remoteOffset, ncclWindow_t localWnd, size_t localOffset, size_t bytes, Coop coop ncclCoopThread{}, DescriptorSmem descriptor ncclGin_None{}, uint32_t optFlags ncclGinOptFlagsDefault, SegmentType bufType ncclGin_SegmentDevice{})get调度与put方向相反的单边传输把远端缓冲区读到本地缓冲区自 NCCL 2.30.3 起可用。参数语义与put对称peer是team内 rank可指向本地回环远端与本地缓冲区用窗口 字节偏移描述bufType默认ncclGin_SegmentDevice。与put不同的是get携带一个optFlags参数默认ncclGinOptFlagsDefault0可选的标志位在 gin_device_common.h 中定义ncclGinOptFlagsMaySkipCreditCheck1 0允许跳过信用检查ncclGinOptFlagsAggregateRequests1 1聚合请求。完成管理flush、flushAsync 与 waitvoid flush(Coop coop, cuda::memory_order ord cuda::memory_order_acquire)flush确保coop中任意线程调度的所有挂起传输操作在本地被消费完毕对put而言意味着源缓冲区可以安全复用它不保证远端 peer 上的完成状态对get而言意味着数据对本地 rank 可见。void flushAsync(ncclTeam team, uint32_t peer, ncclGinRequest_t* request, Coop coop ncclCoopThread{}, uint32_t optFlags ncclGinOptFlagsDefault, DescriptorSmem descriptor ncclGin_None{})flushAsync发起针对单个 peer的非阻塞 flush自 NCCL 2.30.3 起可用。request由调用方提供、由flushAsync初始化调用方随后可用它判断 flush 是否完成void wait(ncclGinRequest_t request, Coop coop ncclCoopThread{}, DescriptorSmem descriptor ncclGin_None{}, cuda::memory_order ord cuda::memory_order_acquire)wait阻塞直到request完成自 NCCL 2.30.3 起可用。这一异步发起 请求对象轮询/等待的模式为计算与通信重叠提供了非阻塞的完成跟踪手段。信号Signals与计数器Counters信号与计数器是 GIN 异步完成语义的两大支柱。信号Signal用于触发远端 peer 上的动作最常见的是在put完成时通知对端计数器Counter用于触发本地 rank 上的动作与信号互补。两者都关联一个可被原子操作的 64 位整数值并采用滚动比较rolling comparison逻辑即使无符号溢出仍保持x x 1的性质该算法在 gin.h 的注释中以rolling_less_equal伪代码给出等待条件为给定值滚动小于等于内部值。强信号与弱信号NCCL 2.30.5 起自 NCCL 2.30.5 起信号分为**强strong与弱weak**两种强信号其可见性意味着同一上下文上、发往同一 peer 的所有此前 put均已落定弱信号只保证其所附着的这一次 put 数据已落定。两类信号各有 Inc/Add 两种对象定义见 gin.h类型结构体成员强 IncncclGin_StrongSignalIncncclGinSignal_t signal强 AddncclGin_StrongSignalAddsignaluint64_t value弱 IncncclGin_WeakSignalIncncclGinSignal_t signal弱 AddncclGin_WeakSignalAddsignaluint64_t value这些对象可作为put、signal等方法的remoteAction参数传入表示收到载荷后把指定signal的值增加若干。其中SignalInc{signalIdx}功能上等价于SignalAdd{signalIdx, 1}但Inc 不能与其他修改信号的操作混用除非中间插入一次信号重置resetSignal。VA 信号任意虚拟地址上的信号除了预分配的信号索引GIN 还支持VA 信号信号位于任意的 VA窗口 偏移对而非预分配的信号槽。对应对象为ncclGin_StrongVASignalInc、ncclGin_StrongVASignalAdd、ncclGin_WeakVASignalInc、ncclGin_WeakVASignalAdd均包含ncclWindow_t signalWindow与size_t signalOffset成员Add 变体另有uint64_t value。使用 VA 信号的前提该窗口必须以NCCL_WIN_STRICT_ORDERING标志注册一旦某地址被用作信号对该地址的所有读写都必须经由 GIN 发起即通过RemoteAction或 GIN 信号方法不能由普通 load/store 访问。自 NCCL 2.30.5 起以下旧信号类型被弃用官方建议显式选用强/弱信号对象ncclGin_VASignalInc、ncclGin_VASignalAdd→ 改用ncclGin_{Strong|Weak}VASignal{Inc|Add}ncclGin_SignalInc、ncclGin_SignalAdd→ 改用ncclGin_{Strong|Weak}Signal{Inc|Add}。对于这些被弃用的信号其强弱由创建设备通信器时ginStrongSignalsRequired字段的值决定该字段在 device_setup.rst 中说明默认true。若将ginStrongSignalsRequired设为false而内核仍使用强信号操作则行为未定义。信号方法ncclGin提供四类信号相关方法void signal(ncclTeam team, int peer, RemoteAction remoteAction, Coop coop, DescriptorSmem descriptor, cuda::thread_scope alreadyReleased, cuda::thread_scope expected_scope)signal实现不带数据传输的显式信号通知参数是put的参数子集——即只发信号、不搬数据。uint64_t readSignal(ncclGinSignal_t signal, int bits 64, cuda::memory_order ord cuda::memory_order_acquire) void waitSignal(Coop coop, ncclGinSignal_t signal, uint64_t least, int bits 64, cuda::memory_order ord cuda::memory_order_acquire) void resetSignal(ncclGinSignal_t signal)readSignal返回signal数值的低bits位waitSignal等待signal数值的低bits位达到或超过leastresetSignal将信号值重置为0注意该方法不能与对信号的其他并发修改竞争。针对 VA 信号还有一组同名重载参数由ncclGinSignal_t换成ncclWindow_t signalWindow, size_t signalOffset。此外gin.h 中还提供了信号影子值shadow相关的高级方法getSignalShadowPtr、increaseSignalShadow、waitSignalMeetShadow、waitSignalFollowShadow每个信号都有专属的影子副本用户可自由操作供复杂的按序消费场景使用。计数器方法计数器使用与信号完全对称的 API但有两个关键差异计数器用于本地动作localAction唯一的动作对象是ncclGin_CounterInc成员为ncclGinCounter_t counter计数器最多只能存储 56 位bits默认56少于uint64_t的完整位宽。uint64_t readCounter(ncclGinCounter_t counter, int bits 56, cuda::memory_order ord cuda::memory_order_acquire) void waitCounter(Coop coop, ncclGinCounter_t counter, uint64_t least, int bits 56, cuda::memory_order ord cuda::memory_order_acquire) void resetCounter(ncclGinCounter_t counter)这三个方法分别对应readSignal/waitSignal/resetSignal的语义读低 bits 位、等待达到阈值、重置为 0。在 gin.h 中ncclGin_CounterInc被标记为弃用建议改用ncclGin_WeakCounterInc弱计数器递增只保证所附着的 put 在本地完成。网络屏障ncclGinBarrierSession跨节点的内核协作需要网络级屏障ncclGinBarrierSession即代表一个网络屏障会话gin_barrier.h 中同时定义了对应的宿主端创建函数ncclGinBarrierCreateRequirement。构造rail 专用变体与通用变体ncclGinBarrierSession(Coop coop, ncclGin gin, ncclTeamTagRail tag, uint32_t index)这是 rail 专用的构造变体coop表示协作线程组ncclCoopThread/ncclCoopWarp/ncclCoopCta见 deviceapi.rst 的 Thread Groups 一节gin是先前初始化的ncclGin对象ncclTeamTagRail表示屏障作用于与本地 rank 同一 rail 的所有 peerindex标识底层使用的屏障槽位——每个coop应当使用不同的 index通常设为blockIdx.x以保证 CTA 之间唯一。ncclGinBarrierSession(Coop coop, ncclGin gin, ncclTeam team, ncclGinBarrierHandle handle, uint32_t index)这是通用变体用于与 world team 等而非 rail team的 rank 通信。它要求显式传入team并额外接受handle参数指明底层屏障的位置通常设为设备通信器的railGinBarrier字段。从 gin_barrier.h 可以看出ncclGinBarrierSession还提供绑定ncclGinAllContextsfence 覆盖所有上下文的构造重载且会话对象不可拷贝Sessions are not copyable并支持带timeoutCycles的sync重载用于超时控制。sync 与 fence 级别void sync(Coop coop, cuda::memory_order order, ncclGinFenceLevel fence ncclGinFenceLevel::Put | ncclGinFenceLevel::Get)sync同步参与屏障会话的所有 team 成员的所有线程。fence是位标志枚举选择屏障返回后必须完成的先前的网络操作省略时默认Put | Get即调用方未显式选择也能获得最强保证。各取值定义在 gin_barrier.h 的ncclGinFenceLevel枚举中取值含义ncclGinFenceLevel::None纯同步不做任何 drain排空ncclGinFenceLevel::Put屏障返回后其他 team 成员在屏障前发出的、以调用方 rank 为目标的 put在调用方内存中可见ncclGinFenceLevel::Get屏障返回后调用方 rank 在屏障前发出的 get 已落入其本地内存ncclGinFenceLevel::Relaxed已弃用的None别名仅为源码级向后兼容保留新代码应使用Nonefence是位标志可通过按位或组合同时请求 Put 与 Get 语义就传ncclGinFenceLevel::Put | ncclGinFenceLevel::Get。在 gin_barrier.h 中可以看到为operator|/operator提供的组合运算符实现。屏障的实现细节见 gin_barrier__funcs.h其中syncInternal会根据fence是否包含Put来决定参与的 peer 信号数量并处理Get的落地检查——这也是默认最强保证在实现层面的体现。实战纯 GIN AlltoAll 内核仓库的 02_alltoall_gin 示例配套说明见 README.md把上述 API 完整串起来所有 peer 数据都走网络正是 GIN 跨节点场景的基线用法。宿主端创建带 GIN 资源的设备通信器#define NCCL_DEVICE_CTA_COUNT 16 #define NCCL_DEVICE_THREADS_PER_CTA 512 ncclDevCommRequirements reqs NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER; reqs.worldGinBarrierCount NCCL_DEVICE_CTA_COUNT; // ncclGinBarrierSessionworld team槽位 reqs.ginSignalCount NCCL_DEVICE_CTA_COUNT; // 每个 CTAblockIdx.x一个信号索引 reqs.ginConnectionType NCCL_GIN_CONNECTION_FULL; // 全连接每个 rank 连到所有 peer NCCLCHECK(ncclDevCommCreate(comm, reqs, devComm));要点与 device_setup.rst 的ncclDevCommRequirements字段一一对应worldGinBarrierCount为ncclGinBarrierSessionworld team 网络屏障预留槽位自 NCCL 2.30.3 起可用ginSignalCount为完成信号预留槽位保证从 id0 开始编号二者都应等于内核发射网格的gridDim.x这里都是 16使每个线程块用blockIdx.x作为自己的屏障索引与信号索引ginConnectionType NCCL_GIN_CONNECTION_FULL建立全连接若无法提供 GINncclDevCommCreate会失败。连接类型枚举NCCL_GIN_CONNECTION_NONE/FULL/RAIL见 device_setup.rst 的ncclGinConnectionType_t。创建前还应通过ncclCommQueryProperties查询props.deviceApiSupport与props.ginType若ginType NCCL_GIN_TYPE_NONE则无法创建 GIN 通信器示例中对此做了显式检查。缓冲区必须用ncclMemAlloc分配并通过ncclCommWindowRegister注册为对称窗口NCCL_WIN_COLL_SYMMETRIC这是 Device API 数据搬移的前提。内核端ncclGin 信号 屏障 puttemplate typename T __global__ void PureGinAlltoAllKernel(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, struct ncclDevComm devComm) { int ginContext 0; // 单上下文 unsigned int signalIndex blockIdx.x; ncclGin gin { devComm, ginContext }; uint64_t signalValue gin.readSignal(signalIndex); ncclGinBarrierSessionncclCoopCta bar { ncclCoopCta(), gin, ncclTeamTagWorld(), blockIdx.x }; bar.sync(ncclCoopCta(), cuda::memory_order_acquire, ncclGinFenceLevel::None); int tid threadIdx.x blockIdx.x * blockDim.x; int nthreads blockDim.x * gridDim.x; const size_t size count * sizeof(T); for (int r tid; r devComm.nRanks; r nthreads) { gin.put(ncclTeamWorld(devComm), r, recvwin, recvoffset devComm.rank * size, sendwin, sendoffset r * size, size, ncclGin_WeakSignalInc{signalIndex}); } int receivingCta (devComm.rank % nthreads) / blockDim.x; if (blockIdx.x receivingCta) gin.waitSignal(ncclCoopCta(), signalIndex, signalValue devComm.nRanks); gin.flush(ncclCoopCta()); bar.sync(ncclCoopCta(), cuda::memory_order_release, ncclGinFenceLevel::None); }逐段解读其与文档 API 的对应关系构造ncclGinncclGin gin { devComm, ginContext }使用上下文 0。生产代码可轮转多个上下文devComm.ginContextCount提升吞吐。入口屏障ncclGinBarrierSessionncclCoopCta以ncclTeamTagWorld()blockIdx.x构造bar.sync在所有 rank 就绪后才放行数据交换此处 fence 显式传ncclGinFenceLevel::None纯同步、不做 drain因为真正的完成保证由后面的信号与 flush 承担。逐 peer 发 put循环内每个线程为不同目的 rank 发一次gin.putremoteAction传ncclGin_WeakSignalInc{signalIndex}——数据落定后目的 rank 上该信号槽位自增 1。AlltoAll 的每个线程索引只需在本 rank 内唯一tid步进nthreads即可随后用该索引选出目的 peer。信号化完成等待所有 peer 发往本 rank 的 put 都会使signalIndex blockIdx.x那个槽位自增一次共需nRanks次。因此只让一个接收 CTAreceivingCta由(devComm.rank % nthreads) / blockDim.x计算执行gin.waitSignal(..., signalValue devComm.nRanks)等待所有入站 put 落定其余 CTA 只发 put 和 flush。flush 收尾gin.flush(ncclCoopCta())保证本地发送缓冲区可安全复用不保证远端完全落定远端完成由waitSignal保证随后再次bar.syncrelease确保所有 rank 完成集体操作后才退出内核。示例还在main.cu中完成了数据初始化my_rank * 1000 dest_rank * 100 element_idx的可验证模式、cudaStreamSynchronize后回拷校验PASSED/FAILED以及ncclDevCommDestroy、ncclCommWindowDeregister、ncclMemFree等清理流程——完整可运行的代码可直接参考该文件。版本兼容与迁移提示若目标 NCCL 版本早于 2.30文档deviceapi.rst 的 Compatibility adjustments 一节给出如下调整建议NCCL 2.28.7 基线GIN 自该版本起可用早于 2.30无worldGinBarrierCountncclGinBarrierSession仅可用于 rail 连通对应字段为railGinBarrierCount若要 world team 的 GIN 屏障将barrierCount设为 CTA 数等于gridDim.x内核中改用混合ncclBarrierSessionncclTeamTagWorld()早于 2.29.7无ginConnectionType将弃用的ginForceEnable设为true以启用全连接等价于NCCL_GIN_CONNECTION_FULLginConnectionType自 NCCL 2.29.7 起可用ginForceEnable自 2.29.7 起弃用。另外注意ncclGin::get、flushAsync、wait自NCCL 2.30.3起可用强/弱信号、VA 信号语义自NCCL 2.30.5起可用GPI 后端自NCCL 2.30.6起作为实验特性加入——编写内核前请核对所用 NCCL 版本与 device_gin.rst 中各 API 标注的引入版本。小结GIN 把网络通信的发起权从 CPU 移交给了 GPU 内核ncclGin封装 GIN 上下文与多后端分发put/get提供单边传输信号/计数器提供远端/本地的异步完成语义含 2.30.5 起的强/弱之分与 VA 信号ncclGinBarrierSession提供带 fence 级别的跨节点屏障。结合仓库中的 API 参考、使用指南 与 完整示例你可以直接以此为基础编写自己的纯 GIN 内核或将其与 LSA、Multimem 结合构建面向多节点、拓扑感知的混合通信方案。赞分享高性能计算通信深度学习分布式训练【免费下载链接】ncclOptimized primitives for collective multi-GPU communication项目地址https://gitcode.com/gh_mirrors/nc/nccl点击查看免费下载相关推荐NCCL Device API 纯 GIN AlltoAll 实战用 GPU 直接发起跨节点网络通信NCCL Device API 纯 GIN AlltoAll 实战用 GPU 直接发起跨节点网络通信 本文以 NCCL 官方示例 02_alltoall_gi高性能计算通信深度学习分布式训练NCCL Device API 纯 GIN AlltoAll 实战从 GPU Kernel 发起全网状集合通信NCCL Device API 纯 GIN AlltoAll 实战从 GPU Kernel 发起全网状集合通信 导读 本文深入解析 NCCL 官方设备 API高性能计算通信深度学习分布式训练NCCL Device API 完全指南从 GPU Kernel 内直接发起通信的 API 体系详解NCCL Device API 完全指南从 GPU Kernel 内直接发起通信的 API 体系详解 NCCL 自 2.28 起提供了 Device API高性能计算通信深度学习分布式训练创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考