NCCL学习笔记-代码解析

最新推荐文章于 2025-04-27 16:56:29 发布
最新推荐文章于 2025-04-27 16:56:29 发布
阅读量877 收藏 21
点赞数 20
分类专栏： NVIDIA NCCL 文章标签： 学习 笔记 算法
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.youkuaiyun.com/qq_51234298/article/details/145001442
版权

Example1:单进程、单线程、多设备（单进程多设备）

单进程下，使用ncclCommInitAll（）

// 声明一个 ncclComm_t 类型的数组 comms，长度为 4，用于存储 NCCL 通信器对象。
ncclComm_t comms[4];
// 声明一个整数数组 devs，包含四个元素，分别为 0、1、2 和 3，
int devs[4] = { 0, 1, 2, 3 };
// 该函数的第一个参数 comms 是一个指向 ncclComm_t 数组的指针，用于存储初始化后的通信器。
// 第二个参数 4 表示要使用的设备数量，即 devs 数组中的元素个数。
// 第三个参数 devs 是一个指向存储设备编号的数组的指针，指定了要使用的设备。
// 该函数会根据提供的设备编号初始化相应的 NCCL 通信器，以便后续进行设备间的通信操作，
// 如集合通信（all-reduce、reduce、broadcast 等）。
ncclCommInitAll(comms, 4, devs);

程序结束可以使用：

for (int i=0; i<4; i++)
  ncclCommDestroy(comms[i]);

完整工作案例（单进程多设备）

#include <stdlib.h>
#include <stdio.h>
#include "cuda_runtime.h"
#include "nccl.h"

// 定义CUDA错误检查宏，用于捕获CUDA API调用中的错误
#define CUDACHECK(cmd) do {                         \
  cudaError_t err = cmd;                            \
  if (err != cudaSuccess) {                         \
    printf("Failed: Cuda error %s:%d '%s'\n",       \
        __FILE__,__LINE__,cudaGetErrorString(err)); \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)

// 定义NCCL错误检查宏，用于捕获NCCL API调用中的错误
#define NCCLCHECK(cmd) do {                         \
  ncclResult_t res = cmd;                           \
  if (res != ncclSuccess) {                         \
    printf("Failed, NCCL error %s:%d '%s'\n",       \
        __FILE__,__LINE__,ncclGetErrorString(res)); \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)

int main(int argc, char* argv[])
{
  ncclComm_t comms[4];  // 定义NCCL通信器数组，用于管理多个设备的通信

  // 管理4个设备
  int nDev = 4;  // 设备数量
  int size = 32*1024*1024;  // 每个设备上缓冲区的大小（32MB）
  int devs[4] = { 0, 1, 2, 3 };  // 设备ID数组

  // 分配和初始化设备缓冲区
  float** sendbuff = (float**)malloc(nDev * sizeof(float*));  // 发送缓冲区
  float** recvbuff = (float**)malloc(nDev * sizeof(float*));  // 接收缓冲区
  cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev);  // CUDA流数组

  // 为每个设备分配内存并初始化
  for (int i = 0; i < nDev; ++i) {
    CUDACHECK(cudaSetDevice(i));  // 设置当前设备
    CUDACHECK(cudaMalloc((void**)sendbuff + i, size * sizeof(float)));  // 分配发送缓冲区
    CUDACHECK(cudaMalloc((void**)recvbuff + i, size * sizeof(float)));  // 分配接收缓冲区
    CUDACHECK(cudaMemset(sendbuff[i], 1, size * sizeof(float)));  // 初始化发送缓冲区为1
    CUDACHECK(cudaMemset(recvbuff[i], 0, size * sizeof(float)));  // 初始化接收缓冲区为0
    CUDACHECK(cudaStreamCreate(s+i));  // 创建CUDA流
  }

  // 初始化NCCL
  NCCLCHECK(ncclCommInitAll(comms, nDev, devs));  // 初始化NCCL通信器

  // 调用NCCL通信API。当每个线程使用多个设备时，需要使用Group API
  NCCLCHECK(ncclGroupStart());  // 开始NCCL通信组
  for (int i = 0; i < nDev; ++i)
    NCCLCHECK(ncclAllReduce((const void*)sendbuff[i], (void*)recvbuff[i], size, ncclFloat, ncclSum,
        comms[i], s[i]));  // 执行AllReduce操作，将发送缓冲区的数据求和并存储到接收缓冲区
  NCCLCHECK(ncclGroupEnd());  // 结束NCCL通信组

  // 同步CUDA流以等待NCCL操作完成
  for (int i = 0; i < nDev; ++i) {
    CUDACHECK(cudaSetDevice(i));  // 设置当前设备
    CUDACHECK(cudaStreamSynchronize(s[i]));  // 同步CUDA流
  }

  // 释放设备缓冲区
  for (int i = 0; i < nDev; ++i) {
    CUDACHECK(cudaSetDevice(i));  // 设置当前设备
    CUDACHECK(cudaFree(sendbuff[i]));  // 释放发送缓冲区
    CUDACHECK(cudaFree(recvbuff[i]));  // 释放接收缓冲区
  }

  // 销毁NCCL通信器
  for(int i = 0; i < nDev; ++i)
      ncclCommDestroy(comms[i]);  // 销毁NCCL通信器

  printf("Success \n");  // 打印成功信息
  return 0;
}

1.CUDA和NCCL初始化：
代码首先初始化了CUDA和NCCL环境，并为每个设备分配了内存和CUDA流。
cudaMalloc用于在设备上分配内存，cudaMemset用于初始化内存。

2.NCCL通信器初始化：
ncclCommInitAll用于初始化NCCL通信器，使得多个设备之间可以进行通信。

3.NCCL AllReduce操作：
ncclAllReduce是NCCL中的一个集合通信操作，它将所有设备的发送缓冲区中的数据求和，并将结果存储到每个设备的接收缓冲区中。
ncclGroupStart和ncclGroupEnd用于将多个NCCL操作组合在一起，确保它们作为一个整体执行。

4.同步和资源释放：
代码通过cudaStreamSynchronize确保所有CUDA流中的操作完成。
最后，代码释放了设备上的内存，并销毁了NCCL通信器。

5.错误处理：
使用CUDACHECK和NCCLCHECK宏来捕获和处理CUDA和NCCL API调用中的错误。

在深度学习训练中，假设有 4 个 GPU，每个 GPU 计算了一部分梯度。使用 ncclAllReduce 可以将所有 GPU 的梯度求和，并将最终的梯度同步到每个 GPU 上。

Example2:一个设备每个进程或线程（多进程单设备）

#include <stdio.h>
#include "cuda_runtime.h"
#include "nccl.h"
#include "mpi.h"
#include <unistd.h>
#include <stdint.h>
#include <stdlib.h>

// MPI 错误检查宏
#define MPICHECK(cmd) do {                          \
  int e = cmd;                                      \
  if( e != MPI_SUCCESS ) {                          \
    printf("Failed: MPI error %s:%d '%d'\n",        \
        __FILE__,__LINE__, e);                      \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)

// CUDA 错误检查宏
#define CUDACHECK(cmd) do {                         \
  cudaError_t e = cmd;                              \
  if( e != cudaSuccess ) {                          \
    printf("Failed: Cuda error %s:%d '%s'\n",       \
        __FILE__,__LINE__,cudaGetErrorString(e));   \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)

// NCCL 错误检查宏
#define NCCLCHECK(cmd) do {                         \
  ncclResult_t r = cmd;                             \
  if (r!= ncclSuccess) {                            \
    printf("Failed, NCCL error %s:%d '%s'\n",       \
        __FILE__,__LINE__,ncclGetErrorString(r));   \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)

// 计算主机名的哈希值
static uint64_t getHostHash(const char* string) {
  uint64_t result = 5381;
  for (int c = 0; string[c] != '\0'; c++){
    result = ((result << 5) + result) ^ string[c];
  }
  return result;
}

// 获取主机名
static void getHostName(char* hostname, int maxlen) {
  gethostname(hostname, maxlen);
  for (int i=0; i< maxlen; i++) {
    if (hostname[i] == '.') {
        hostname[i] = '\0';
        return;
    }
  }
}

int main(int argc, char* argv[]) {
  int size = 32*1024*1024;  // 缓冲区大小（32MB）
  int myRank, nRanks, localRank = 0;  // MPI 进程信息

  // 初始化 MPI
  MPICHECK(MPI_Init(&argc, &argv));
  MPICHECK(MPI_Comm_rank(MPI_COMM_WORLD, &myRank));
  MPICHECK(MPI_Comm_size(MPI_COMM_WORLD, &nRanks));

  // 计算本地 Rank
  uint64_t hostHashs[nRanks];
  char hostname[1024];
  getHostName(hostname, 1024);  // 获取主机名
  hostHashs[myRank] = getHostHash(hostname);  // 计算主机名的哈希值
  MPICHECK(MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, hostHashs, sizeof(uint64_t), MPI_BYTE, MPI_COMM_WORLD));  // 收集所有进程的主机名哈希值
  for (int p=0; p<nRanks; p++) {
    if (p == myRank) break;
    if (hostHashs[p] == hostHashs[myRank]) localRank++;  // 计算本地 rank
  }

  // 初始化 NCCL
  ncclUniqueId id;
  ncclComm_t comm;
  if (myRank == 0) ncclGetUniqueId(&id);  // 由 rank 0 生成 NCCL 唯一 ID
  MPICHECK(MPI_Bcast((void *)&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD));  // 广播 NCCL 唯一 ID

  // 选择 GPU 并分配设备内存
  float *sendbuff, *recvbuff;
  cudaStream_t s;
  CUDACHECK(cudaSetDevice(localRank));  // 根据本地 rank 选择 GPU
  CUDACHECK(cudaMalloc(&sendbuff, size * sizeof(float)));  // 分配发送缓冲区
  CUDACHECK(cudaMalloc(&recvbuff, size * sizeof(float)));  // 分配接收缓冲区
  CUDACHECK(cudaStreamCreate(&s));  // 创建 CUDA 流

  // 初始化 NCCL 通信器
  NCCLCHECK(ncclCommInitRank(&comm, nRanks, id, myRank));

  // 执行 NCCL 集合通信
  NCCLCHECK(ncclAllReduce((const void*)sendbuff, (void*)recvbuff, size, ncclFloat, ncclSum, comm, s));  // 执行 AllReduce 操作
  CUDACHECK(cudaStreamSynchronize(s));  // 同步 CUDA 流

  // 释放资源
  CUDACHECK(cudaFree(sendbuff));  // 释放发送缓冲区
  CUDACHECK(cudaFree(recvbuff));  // 释放接收缓冲区
  ncclCommDestroy(comm);  // 销毁 NCCL 通信器

  // 结束 MPI
  MPICHECK(MPI_Finalize());

  printf("[MPI Rank %d] Success \n", myRank);
  return 0;
}

与单进程多设备模式的区别
  1. 进程与设备的关系
    单进程多设备：
    一个进程管理多个 GPU 设备。
    使用 ncclCommInitAll 初始化 NCCL 通信器。
    适用于单节点多 GPU 场景。
    多进程单设备（当前代码）：
    每个进程管理一个 GPU 设备。
    使用 ncclCommInitRank 初始化 NCCL 通信器。
    适用于多节点多 GPU 场景。
  2. 设备选择
    单进程多设备：
    在代码中显式指定每个设备（如 cudaSetDevice(i)）。
    设备选择由进程内部逻辑决定。
    多进程单设备：
    根据 localRank 选择设备（如 cudaSetDevice(localRank)）。
    设备选择基于主机名和 MPI rank。
  3. 通信器初始化
    单进程多设备：
    使用 ncclCommInitAll 一次性初始化所有设备的通信器。
    多进程单设备：
    使用 ncclCommInitRank 为每个进程初始化通信器。
    需要 MPI 广播 NCCL 唯一 ID。
  4. 适用场景
    单进程多设备：
    适用于单节点多 GPU 训练。
    例如，一台服务器上有 8 个 GPU，一个进程管理所有 GPU。
    多进程单设备：
    适用于多节点多 GPU 训练。
    例如，多个节点，每个节点有 4 个 GPU，每个 GPU 由一个 MPI 进程管理。

Example3:多设备在一个线程（多进程多设备）

#include <stdio.h>
#include "cuda_runtime.h"
#include "nccl.h"
#include "mpi.h"
#include <unistd.h>
#include <stdint.h>

// MPI 错误检查宏
#define MPICHECK(cmd) do {                          \
  int e = cmd;                                      \
  if( e != MPI_SUCCESS ) {                          \
    printf("Failed: MPI error %s:%d '%d'\n",        \
        __FILE__,__LINE__, e);                      \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)

// CUDA 错误检查宏
#define CUDACHECK(cmd) do {                         \
  cudaError_t e = cmd;                              \
  if( e != cudaSuccess ) {                          \
    printf("Failed: Cuda error %s:%d '%s'\n",       \
        __FILE__,__LINE__,cudaGetErrorString(e));   \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)

// NCCL 错误检查宏
#define NCCLCHECK(cmd) do {                         \
  ncclResult_t r = cmd;                             \
  if (r!= ncclSuccess) {                            \
    printf("Failed, NCCL error %s:%d '%s'\n",       \
        __FILE__,__LINE__,ncclGetErrorString(r));   \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)

// 计算主机名的哈希值
static uint64_t getHostHash(const char* string) {
  uint64_t result = 5381;
  for (int c = 0; string[c] != '\0'; c++){
    result = ((result << 5) + result) ^ string[c];
  }
  return result;
}

// 获取主机名
static void getHostName(char* hostname, int maxlen) {
  gethostname(hostname, maxlen);
  for (int i=0; i< maxlen; i++) {
    if (hostname[i] == '.') {
        hostname[i] = '\0';
        return;
    }
  }
}

int main(int argc, char* argv[]) {
  int size = 32*1024*1024;  // 缓冲区大小（32MB）

  int myRank, nRanks, localRank = 0;  // MPI 进程信息

  // 初始化 MPI
  MPICHECK(MPI_Init(&argc, &argv));
  MPICHECK(MPI_Comm_rank(MPI_COMM_WORLD, &myRank));  // 获取当前进程的 rank
  MPICHECK(MPI_Comm_size(MPI_COMM_WORLD, &nRanks));  // 获取总的进程数量

  // 计算本地 Rank
  uint64_t hostHashs[nRanks];
  char hostname[1024];
  getHostName(hostname, 1024);  // 获取主机名
  hostHashs[myRank] = getHostHash(hostname);  // 计算主机名的哈希值
  MPICHECK(MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, hostHashs, sizeof(uint64_t), MPI_BYTE, MPI_COMM_WORLD));  // 收集所有进程的主机名哈希值
  for (int p=0; p<nRanks; p++) {
     if (p == myRank) break;
     if (hostHashs[p] == hostHashs[myRank]) localRank++;  // 计算本地 rank
  }

  // 每个进程使用两个 GPU
  int nDev = 2;  // 每个进程管理的 GPU 数量

  // 分配设备缓冲区
  float** sendbuff = (float**)malloc(nDev * sizeof(float*));  // 发送缓冲区
  float** recvbuff = (float**)malloc(nDev * sizeof(float*));  // 接收缓冲区
  cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev);  // CUDA 流

  // 根据本地 rank 选择 GPU 并分配设备内存
  for (int i = 0; i < nDev; ++i) {
    CUDACHECK(cudaSetDevice(localRank*nDev + i));  // 设置当前 GPU
    CUDACHECK(cudaMalloc(sendbuff + i, size * sizeof(float)));  // 分配发送缓冲区
    CUDACHECK(cudaMalloc(recvbuff + i, size * sizeof(float)));  // 分配接收缓冲区
    CUDACHECK(cudaMemset(sendbuff[i], 1, size * sizeof(float)));  // 初始化发送缓冲区为 1
    CUDACHECK(cudaMemset(recvbuff[i], 0, size * sizeof(float)));  // 初始化接收缓冲区为 0
    CUDACHECK(cudaStreamCreate(s+i));  // 创建 CUDA 流
  }

  // 初始化 NCCL
  ncclUniqueId id;
  ncclComm_t comms[nDev];

  // 生成 NCCL 唯一 ID 并广播给所有进程
  if (myRank == 0) ncclGetUniqueId(&id);  // 由 rank 0 生成 NCCL 唯一 ID
  MPICHECK(MPI_Bcast((void *)&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD));  // 广播 NCCL 唯一 ID

  // 初始化 NCCL 通信器，使用 Group API
  NCCLCHECK(ncclGroupStart());
  for (int i=0; i<nDev; i++) {
     CUDACHECK(cudaSetDevice(localRank*nDev + i));  // 设置当前 GPU
     NCCLCHECK(ncclCommInitRank(comms+i, nRanks*nDev, id, myRank*nDev + i));  // 初始化 NCCL 通信器
  }
  NCCLCHECK(ncclGroupEnd());

  // 使用 NCCL 通信 API，使用 Group API
  NCCLCHECK(ncclGroupStart());
  for (int i=0; i<nDev; i++)
     NCCLCHECK(ncclAllReduce((const void*)sendbuff[i], (void*)recvbuff[i], size, ncclFloat, ncclSum,
           comms[i], s[i]));  // 执行 AllReduce 操作
  NCCLCHECK(ncclGroupEnd());

  // 同步 CUDA 流以完成 NCCL 通信
  for (int i=0; i<nDev; i++)
      CUDACHECK(cudaStreamSynchronize(s[i]));  // 同步 CUDA 流

  // 释放设备内存
  for (int i=0; i<nDev; i++) {
     CUDACHECK(cudaFree(sendbuff[i]));  // 释放发送缓冲区
     CUDACHECK(cudaFree(recvbuff[i]));  // 释放接收缓冲区
  }

  // 销毁 NCCL 通信器
  for (int i=0; i<nDev; i++) {
     ncclCommDestroy(comms[i]);  // 销毁 NCCL 通信器
  }

  // 结束 MPI
  MPICHECK(MPI_Finalize());

  printf("[MPI Rank %d] Success \n", myRank);
  return 0;
}

总结上述的：
在这里插入图片描述

open-mmlab学习笔记
xiaomu_347的博客
05-15 261
OpenMMLab 是一个由香港中文大学多媒体实验室 (MMLab) 发起的开源项目，旨在建立一个通用的、模块化的深度学习模型框架，以便于研究者和开发者能够快速构建和部署深度学习模型。这些框架在计算机视觉领域的研究和应用中发挥了重要作用，帮助研究者和开发者快速实现和部署深度学习模型。: 一个全面的、模块化的、灵活的目标检测工具箱，支持单阶段和双阶段检测器。: 用于语义分割、实例分割和全景分割的框架，提供多种分割网络结构。: 用于生成模型的框架，如条件图像生成、文本到图像生成等。
LLM学习笔记
weixin_40103562的博客
10-16 839
常见术语。
NCCL源码解析④：建图过程
OneFlow深度学习框架
06-27 3415
上次分析了NCCL对机器PCI系统进行拓扑分析的过程，产出的结果为xml格式，接下来，NCCL会根据这个xml进图的建立过程以便之后进行路径搜索。ncclTopoGetSystem的最后会执行ncclTopoGetSystemFromXml将xml格式转成图格式。ncclResult_t ncclTopoGetSystemFromXml(struct ncc...
NCCL 源码详解总目录
这个地瓜真是甜啊
02-05 6023
NCCL 源码解析总目录
【大模型】通信元语和相关概念|NCCL梯度|Allreduce|Scatter|Broadcast|Gather
bandaoyu的note
03-24 1540
ReduceReduce（归约） 是一种聚合操作，与“拆分”相反，它的核心逻辑是将多个数据元素合并成一个结果。拆分是 Map 的职责，而非 Reduce。Reduce 的“归约”本质是通过迭代合并，逐步缩小数据规模，最终得到全局结果。以下是关键要点：1.定义与功能Reduce通过一个二元操作（如加法、取最大值等）逐步合并数据，最终输出一个单一值。例如：对执行 Reduce（求和），结果为10。2.与 Map 的关系在MapReduce模型中：Map。
NCCL源码详解3：通信器初始化ncclCommInitRank() 含视频教程
爱串门的小马驹博客
05-24 3080
NCCL源码解析解读分析，ncclCommInitRank()根据rank号和CUDA设备uniqueID等，完成NCCL通信器初始化。ncclCommInitRankDev()使用 ncclAsyncLaunch 异步启动 ncclCommInitRankFunc 函数来初始化通信。 异步启动通信器初始化ncclCommInitRankFunc()
NCCL test源码分析
weixin_48329334的博客
09-14 1907
Nccl test有几个测试文件，每一个测试文件会单独生成一个测试程序，并以"_perf"为后缀生成可执行的文件。每个测试源码文件会和两个公共文件common.cu和timer.cc，以及nccl、cuda、mpi这三个lib一起编译生成可执行的文件。
NCCL源码解读5：拓扑识别感知整体思路总览
爱串门的小马驹博客
06-17 1868
依据物理拓扑结构，来进行通信，可达到性能更优，这也是Nvidia NCCL的核心功能。NCCL拓扑识别的整体思路： 1、物理拓扑构建 2、通信路径计算（每个GPU/网卡到其它GPU，网卡的最优路径。） 3、逻辑拓扑构建(通信通道检索) 哈哈哈，整体思路是不是特别简单。就是先获取物理拓扑图，然后计算通信路径(方便逻辑拓扑构建)，根据通信路径构建逻辑拓扑，例如ring，tree逻辑拓扑，指明哪个GPU和哪个GPU通信。
理解NCCL源码必看：详细论述NCCL源码做了哪些事
这个地瓜真是甜啊
03-21 7470
详解NCCL 源码的实现的功能以及如何实现的
【分布式】入门级NCCL多机并行实践 - 02
2401_84433974的博客
04-27 2306
从上图可知，1.6版本左右的PyTorch调用结构如下，最后在ProcessGroup.hpp可以找到对NCCL、Gloo和MPI的调用。而这些蓝色的部分就是基本的分布式通讯库，他们负责实现通讯和一部分计算功能。
【吴恩达机器学习课程深度体验】：Pytorch项目实战技巧与心得分享（珍贵学习笔记）
[【吴恩达机器学习课程深度体验】：Pytorch项目实战技巧与心得分享（珍贵学习笔记）](https://cdn.educba.com/academy/wp-content/uploads/2021/10/PyTorch-Model.jpg) # 摘要 本文系统地介绍了吴恩达机器学习课程...
NCCL-用于集体多GPU通信的优化原语-C/C++开发
05-26
NCCL-集体多GPU通讯的优化原语NCCL集体多GPU通讯的优化原语。 简介NCCL（发音为“镍”）是用于GPU的标准集体通信例程的独立库，可实现全缩减，全收集，缩减，广播和缩减分散。 它已经过优化，可以在使用PCIe，NVLink，NVswitch的平台以及使用InfiniBand Verbs或TCP / IP套接字的网络上实现高带宽。 NCCL支持在单个节点上或跨节点安装的任意数量的GPU
DeepSpeed MoE 加速引擎实战指南：超大专家模型训练与推理全流程解析
最新发布
努力分享一些人工智能相关的知识干货！
04-27 705
随着参数规模迈向千亿甚至万亿级别，传统稠密模型（Dense Model）已难以在硬件资源、训练成本与推理效率之间取得平衡。 稀疏激活的专家模型（Mixture of Experts，MoE）体系应运而生，通过在每次前向计算中仅激活部分子网络（专家），实现了参数量巨大但计算开销可控的突破。 DeepSpeed 作为业界领先的大模型训练与推理加速框架，提供了完整、成熟的 MoE 支持，包括 MoE层实现、稀疏路由调度、数据并行与专家并行融合等高级优化技术。 本篇作为 DeepSpeed 系列第八篇，
NCCL源码解析: 共享内存连接的建立
这个地瓜真是甜啊
02-06 2403
NCCL 源码中两个设备通过共享内存通信初始化
NCCL源码解析⑤：路径计算
OneFlow深度学习框架
07-20 4460
上节NCCL完成了对机器PCI系统拓扑的建图，其中建好的图如下所示，其中GPU之间是通过NVLink连接起来的。为了方便之后的搜索channel，接下来NCCL会先计算GPU和NIC节点到其他任意节点之间的最优路径，以及对应的带宽，即最优路径上所有边的带宽的最小值。那么抽象一下，这个问题可以建模为给定一个无向图，每条边有一个权值，给定查询(u, v)，求节点...
NVIDIA NCCL 源码学习（四）- 建图过程
KIDGIN7439的专栏
10-26 8113
上次分析到nccl对机器PCI系统进行拓扑分析的过程，产出的结果为xml格式，接下来，nccl会根据这个xml进图的建立过程以便之后进行路径搜索。
nccl 源码分析 从 ncclAllReduce 的执行开始认识nccl源代码
eloudy的专栏
01-04 2389
这个info是在ncclAllReduce()中定义的 struct ncclInfo info，其中info->comm 是ncclAllReduce 传递进来的第五个参数 ncclComm* comm。所以，启动相关的阿cuda kernel等，应该是在后面的 ncclGroupEnd() 中，通过解析 comm->tasks的数据元素来启动的。那么需要看一下 ncclIntruQueueEnqueue 到底对t中的 t->op做了什么解析，t->op是这个函数的第二个参数的op成员；
NVIDIA NCCL 源码学习（十二）- double binary tree
KIDGIN7439的专栏
12-22 7850
上节我们以ring allreduce为例看到了集合通信的过程，但是随着训练任务中使用的gpu个数的扩展，ring allreduce的延迟会线性增长，为了解决这个问题，NCCL引入了tree算法，即double binary tree。
NCCL源码详解4：bootstrapInit()引导网络bootstrap网络连接建立 视频教程
爱串门的小马驹博客
06-04 2246
Nvidia NCCL bootstrapInit()利用已知的根节(rank0)网络地址(从而获得ncclUniqueId中包含的rank0的网络地址)，建立一个环形网络，allgather获取所有rank的信息。
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

MatsumotoChrikk

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值