NCCl API ncclCommInitAll功能实现，测试和结果验证

admin管理员组

文章数量:1442779

NCCl API ncclCommInitAll功能实现，测试和结果验证

测试程序使用的cuda版本，nccl版本和硬件

NCCL 版本 ： 2.19.3

CUDA版本： cuda_12.4.r12.4

硬件 ： 1CPU+2GPU（3090）

ncclCommInitAll API 源代码解析

我们将ncclCommInitAll函数分为6个主要子模块：NVTX功能相关、libcuda.so动态库相关、ncclGetUniqueId函数相关、ncclGroupStart函数相关、ncclCommInitRankDev函数相关和ncclGroupEnd函数相关。如下图所示，每个子模块都会做相应的分析。

NVTX功能相关

我们暂时不关注调式事项，暂时不深入此功能。

libcuda.so动态库相关

N家GPU对应libcuda.so动态库，其他GPU厂家类似动态库可以同名字覆盖此文件，也可以改为其他名字与此动态库共存。

统计此动态库一共使用的cuda driver API如下：cuInit、cuDriverGetVersion、cuGetProcAddress、cuGetErrorString、cuGetErrorName、cuDeviceGet、cuDeviceGetAttribute、cuMemGetAddressRange、cuCtxCreate、cuCtxDestroy、cuCtxGetCurrent、cuCtxSetCurrent、cuCtxGetDevice、cuMemAddressReserve、cuMemAddressFree、cuMemCreate、cuMemGetAllocationGranularity、cuMemExportToShareableHandle、cuMemImportFromShareableHandle、cuMemMap、cuMemRelease、cuMemRetainAllocationHandle、cuMemSetAccess、cuMemUnmap、cuPointerGetAttribute、cuMemGetHandleForAddressRange、cuMulticastAddDevice、cuMulticastBindMem、cuMulticastBindAddr、cuMulticastCreate、cuMulticastGetGranularity和cuMulticastUnbind。并且调用时给出了此API在哪些版本支持，这里不在列举，因为类CUDA Driver API支持与否不在于CUDA版本，而在于类GPU芯片本身是否支持此API功能。这里是不详细分析各个API实现原理和使用方法。

ncclGetUniqueId函数相关

这个函数分析见 ，需要注意ID前后两部分来源和实现原理。

ncclGroupStart函数相关

核心代码是ncclGroupDepth++;不用太多解释。

ncclCommInitRankDev函数相关

ncclCommInitRankDev函数是ncclCommInitAll API功能实现的核心函数，下面会重点分析初始化功能和ncclCommInitRankFunc子模块，其他部分见下图所示：

初始化功能

cudaSetDevice+cudaDeviceGetAttribute（cudaDevAttrComputeCapabilityMajor+cudaDevAttrComputeCapabilityMinor）获取cudaArch+ncclInitKernelsForDevice获取maxLocalSizeBytes+cudaDeviceSetLimit函数设置cudaLimitStackSize属性，相对不复杂，不在详细解析。

bootstrapInit函数主要功能是通过socket+accept+connect+send/recv建立TCP连接并通讯实现数据交互，然后进一步创建代理服务。（后面会用到，与非主线分支，暂时不处理）

建立拓扑

initTransportsRank函数处理逻辑很多，为了避免陷入细节，我们这里只关注下面三个要点，其他部分咱们不分析

获取设备信息：获取当前 GPU 设备的硬件信息，如设备 ID、PCIe 总线 ID、NVLink 连接状态等。

填充拓扑信息：将收集到的设备信息填充到 ncclPeerInfo 结构中，用于后续的拓扑分析和通信优化。

支持多 GPU 架构：适应多种 GPU 连接方式（如 PCIe、NVLink、NVSwitch 等），确保在不同拓扑结构下的性能优化。

##ncclGroupEnd函数相关

所有的NCCL库的APi函数都会调用此函数，对比ncclBroadcast等实际功能API对于ncclCommInitAll这个API的处理与前面都不一样。它走的分支是asyncJobsMain，其他API走的是HeadMain分支。如下图所示，处理逻辑相对比较简单，直接看图就可以抓住核心代码流程。

ncclCommInitAll API调用例子源代码

#include <stdlib.h>
#include <stdio.h>
#include <stdint.h>

#include <sys/socket.h>
#include <arpa/inet.h>
#include <netinet/tcp.h>
#include <netdb.h>
#include <fcntl.h>
#include <poll.h>

#include "cuda_runtime.h"
#include "nccl.h"

#define DEFAULT_ARGC      (2)
char get_gpu_number_cmd[] = "nvidia-smi topo -m";

#if 1
#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)

#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)
#endif

void show_h_A_value(int *h_A, int cnt)
{
    for (int i = 0; i < cnt; ++i) {
        printf("%d ", (int)h_A[i]);
    }
    printf("\n");
}

int main(int argc, char* argv[])
{
  ncclComm_t comms[4];

  //managing 4 devices
  int nDev = 1;
  int devs[4] = { 0, 1, 2, 3 };

  printf("=====================start==============\n");
  printf("[%s +%d %s] ncclCommInitAll/ncclCommDestroy API test example\n", __FILE__, __LINE__, __func__);

  if(DEFAULT_ARGC != argc)
  {
    printf("[%s +%d %s] argc=%d !=%d so exit(0)\n", __FILE__, __LINE__, __func__, argc, DEFAULT_ARGC);
    printf("[%s +%d %s] example ./nccl_CommInitAll_CommDestroy.elf num is the number can used \n", __FILE__, __LINE__, __func__);
    printf("./nccl_CommInitAll_CommDestroy.elf 1 \n");
    printf("./nccl_CommInitAll_CommDestroy.elf 2 \n");
    printf("./nccl_CommInitAll_CommDestroy.elf 4 \n");
	  system(get_gpu_number_cmd);
    exit(0);
  }
  else
  {
    nDev = atoi(argv[1]);
    printf("[%s +%d %s] input nDev=%d\n", __FILE__, __LINE__, __func__, nDev);
  }


  printf("[%s +%d %s] call ncclCommInitAll\n", __FILE__, __LINE__, __func__);
  //initializing NCCL
  NCCLCHECK(ncclCommInitAll(comms, nDev, devs));

  #if 0
  // get some config from comm
  int com_count = 0;
  int com_devid = 0;
  int com_rank = 0;
  for(int i = 0; i < nDev; ++i)
  {
    printf("[%s +%d %s] call ncclCommCount\n", __FILE__, __LINE__, __func__);
    NCCLCHECK(ncclCommCount(comms[i], &com_count));
    printf("[%s +%d %s] [%d] com_count=%d\n", __FILE__, __LINE__, __func__, i, com_count);
    
    printf("[%s +%d %s] call ncclCommCuDevice\n", __FILE__, __LINE__, __func__);
    NCCLCHECK(ncclCommCuDevice(comms[i], &com_devid));
    printf("[%s +%d %s] [%d] com_devid=%d\n", __FILE__, __LINE__, __func__, i, com_devid);
    
    printf("[%s +%d %s] call ncclCommUserRank\n", __FILE__, __LINE__, __func__);
    NCCLCHECK(ncclCommUserRank(comms[i], &com_rank));
    printf("[%s +%d %s] [%d] com_rank=%d\n", __FILE__, __LINE__, __func__, i, com_rank);
  }
  #endif
  //finalizing NCCL
  for(int i = 0; i < nDev; ++i)
  {
    ncclCommDestroy(comms[i]);
  }

  printf("=====================end==============\n");
  return 0;
}

ncclCommInitAll API 运行结果

本文标签： NCCl API ncclCommInitAll功能实现，测试和结果验证