NCCL源码解析: 共享内存连接的建立

已于 2024-03-20 08:46:23 修改
阅读量1.5k 收藏 19
点赞数 21
分类专栏： NCCL 文章标签： NCCL SHM nvidia P2P
于 2024-02-06 10:44:13 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/u014443578/article/details/136053107
版权

文章目录

前言

NCCL 源码解析总目录

我尽量在每个函数之前介绍每个函数的作用，建议先不要投入到函数内部实现，先把函数作用搞清楚，有了整体框架，再回归到细节。

习惯： 我的笔记习惯：为了便于快速理解，函数调用关系通过缩进表示，也可能是函数展开，根据情况而定。

如下

// 调用 proxyConnInit
NCCLCHECK(proxyConnInit(peer, connectionPool, proxyState, (ncclProxyInitReq*) op->reqBuff, (ncclProxyInitResp*) op->respBuff, &op->connection));
// 对函数 proxyConnInit 进行展开，可方便看参数
static ncclResult_t proxyConnInit(struct ncclProxyLocalPeer* peer, struct ncclProxyConnectionPool* connectionPool, struct ncclProxyState* proxyState, ncclProxyInitReq* req, ncclProxyInitResp* resp, struct

如有问题，请留言指正。

图后面再补；
有些遗漏之处，还没涉及，后面补；
闲话后面再补。

概括

recvpeer 表示本卡作为接收端的对端
sendpeer 表示本卡作为发送端的对端

对于每个 channel ，卡与卡之间要建立通信，先通过调用 selectTransport<0>() 建立接收通道，0 表示与 recvpeer 建立通信，再通过selectTransport<1>() 建立发送通道，1表示与 sendpeer 建立通信。
建立通道时会遍历 NTRANSPORTS 4种情况：P2P、共享内存、网络、collNet(collective Network, 还没看，不了解)

struct ncclTransport* ncclTransports[NTRANSPORTS] = {
  &p2pTransport,
  &shmTransport,
  &netTransport,
  &collNetTransport
};

本文重点关注 shmTransport。

接口如下：


struct ncclTransport shmTransport = {
  "SHM",
  shmCanConnect,
  { shmSendSetup, shmSendConnect, shmSendFree, NULL, NULL, NULL, NULL, NULL },
  { shmRecvSetup, shmRecvConnect, shmRecvFree, NULL, NULL, NULL, NULL, NULL }
};

发送建立流程为 shmCanConnect() -> shmSendSetup()
接收建立流程为 shmCanConnect() -> shmRecvSetup()

共享内存相对比较简单，就是在两个 GPU 设备不能进行 nvlink 或者通过switch 进行 P2P 的时候，在系统内存中分配一段空间，两张卡都通过操作这一段共享内存进行数据通信。

详解

1. shmCanConnect()

检查能不能使用共享内存：

  1. 网络快的话，用网络；
  2. 两张卡在一个主机上，要不然看不到同一块内存；
  3. 两张卡都能通过访问 /dev/shm 文件共享内存；
selectTransport<0>(comm, graph, recvData[i]+recvChannels++, c, recvPeer, connIndex, &type)
	static ncclResult_t selectTransport(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclConnect* connect, int channelId, int peer, int connIndex, int* transportType)
	{
		struct ncclPeerInfo* myInfo = comm->peerInfo+comm->rank;
  		struct ncclPeerInfo* peerInfo = comm->peerInfo+peer;
		struct ncclConnector* connector = (type == 1) ? comm->channels[channelId].peers[peer]->send + connIndex :
                                                  comm->channels[channelId].peers[peer]->recv + connIndex;
		NCCLCHECK(transport->canConnect(&ret, comm->topo, graph, myInfo, peerInfo));
		connector->transportComm = transportComm;
		NCCLCHECK(transportComm->setup(comm, graph, myInfo, peerInfo, connect, connector, channelId, connIndex));
	}


// 第一步，检查能不能用共享能存
NCCLCHECK(transport->canConnect(&ret, comm->topo, graph, myInfo, peerInfo));
static ncclResult_t shmCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) 
{
	// 环境变量不能禁止 NCCL_SHM_DISABLE
	if (ncclParamShmDisable() == 1) 
		return ncclSuccess;
	// 检查用网卡是不是速度更快
	// 比较 GPU1 到网卡的 bw 与 GPU1 到 GPU2 的 bw
	// 比较 GPU2 到网卡的 bw 与 GPU1 到 GPU2 的 bw
	// 如果两个都快，那么用网络
	int useNet = 0;
  	NCCLCHECK(ncclTopoCheckNet(topo, info1->busId, info2->busId, &useNet));
  	if (useNet) 
		return ncclSuccess;

	// 还要保证两张卡再同一个主机上
	if (info1->hostHash != info2->hostHash) 
		return ncclSuccess;

	// 确保两张卡的环境能访问同一块内存 /dev/shm
	// info->shmDev = statbuf.st_dev;
	// shmDev 保存的是设备文件的主次设备号，根据这个信息可以决定容器环境中是否可以使用共享内存
	if (info1->shmDev != info2->shmDev) 
		return ncclSuccess;

	// 以上都没问题，就可以用共享内存了
	*ret = 1;
}

2. shmSendSetup

发送方向设置：

  1. 在 /dev/shm 中创建文件，就是在内存上分配内存；
  2. 通过 mmap 映射文件，就是映射内存，返回的首地址 hptr，线程可以通过 hptr 读写此段内存；
  3. 通过 CUDA API 建立映射 hptr ，返回地址 dptr，设备可以通过 dptr 访问该地址；
  4. 记录信息保存在struct shmSendResources* resources->hostHandle 中，资源信息 resources
// 发送设置
// connect : 连接信息缓冲区首地址
// connector : 发送或者接收软件抽象
NCCLCHECK(transportComm->setup(comm, graph, myInfo, peerInfo, connect, connector, channelId, connIndex));
static ncclResult_t shmSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclPeerInfo* myInfo, struct ncclPeerInfo* peerInfo, struct ncclConnect* connectInfo, struct ncclConnector* send, int channelId, int connIndex) 
{
	// 发送缓冲区申请内存
	struct shmSendResources* resources;
	NCCLCHECK(ncclCalloc(&resources, 1));
	send->transportResources = resources;
	// 连接缓冲区，通过强制转换，填充不同类型的连接信息
	struct shmConnectInfo* info = (struct shmConnectInfo*)connectInfo;

	int shmSize = sizeof(struct ncclSendMem);
	info->shmSize = resources->shmSize = shmSize;
	// /dev/shm 可以认为是内存设备的实体文件，对此目录的操作会落到内存的读写上
	// 在该目录创建的文件保存在内存中，大小是由限制，最大内存的一半
	// resources->hostMem 保存内存首地址
	// resources->devHostMem 保存设备要访问共享内存时使用的首地址
	// resources->hostHandle 保存共享内存所有信息
	NCCLCHECK(ncclShmOpen(shmPath, resources->shmSize, (void**)&resources->hostMem, (void**)&resources->devHostMem, 1, &resources->hostHandle));
	ncclResult_t ncclShmOpen(char* shmPath, size_t shmSize, void** shmPtr, void** devShmPtr, int refcount, ncclShmHandle_t* handle)
	{
		// 多申请 4 个字节，要保存 refcount 信息， 引用计数
		const size_t refSize = sizeof(int); /* extra sizeof(int) bytes for reference count */
		const size_t realShmSize = shmSize + refSize;
		// 打开文件，申请内存
		SYSCHECKGOTO(fd = open(shmPath, O_CREAT | O_RDWR, S_IRUSR | S_IWUSR), ret, fail);
		(ftruncate(fd, realShmSize) != 0)
		// 内存申请成功之后，使用 mmap 进行映射，程序操作映射返回的首地址 hptr，就是操作申请的这段空间的内存
		hptr = (char*)mmap(NULL, realShmSize, PROT_READ | PROT_WRITE, MAP_SHARED, fd, 0);
		// 最后 4个字节记录引用计数，新创建的为 1
		*(int*)(hptr + shmSize) = refcount;
		if (devShmPtr) {
			// cudaHostRegister() 把可分页内存标记为锁页内存, 内存不可换出
			// cudaHostRegisterMapped 表示把锁页内存地址映射到设备地址空间
			// 这样，这块存储会有两个地址：一个是从cudaHostAlloc() 或 malloc() 返回的在主机内存地址空间上
			// 另一个在设备存储器上，可以通过cudaHostGetDevicePointer() 取得, 核函数可以通过这个地址访问这段空间
			CUDACHECKGOTO(cudaHostRegister((void*)hptr, realShmSize, cudaHostRegisterMapped), ret, fail);
			// GPU 内部使用 dptr 访问这段空间
			CUDACHECKGOTO(cudaHostGetDevicePointer(&dptr, (void*)hptr, 0), ret, fail);
		}
		// 所有信息保存在 tmphandle 中
		shmHandleInit(fd, shmPath, shmSize, realShmSize, hptr, dptr, create, tmphandle);
		static void shmHandleInit(int fd, char* shmPath, size_t shmSize, size_t realShmSize, char* hptr, void* dptr, bool create, struct shmHandleInternal* handle) 
		{
			handle->fd = fd;
			handle->shmPtr = hptr;
			handle->devShmPtr = dptr;
			handle->shmSize = shmSize;
			handle->realShmSize = realShmSize;
			handle->refcount = (hptr != NULL) ? (int*)(hptr + shmSize) : NULL;
			if (create) {
				int slen = strlen(shmPath);
				handle->shmPath = (char*)malloc(slen + 1);
				memcpy(handle->shmPath, shmPath, slen + 1);
				if (hptr) memset(hptr, 0, shmSize);
			} else {
				handle->shmPath = NULL;
			}
		}
		*shmPtr = hptr;
		if (devShmPtr) 
			*devShmPtr = dptr;
  		*handle = (ncclShmHandle_t)tmphandle;
	}
}

3. shmRecvSetup

接收方向设置：，与发送一样：

  1. 在 /dev/shm 中创建文件，就是在内存上分配内存；
  2. 通过 mmap 映射文件，就是映射内存，返回的首地址 hptr，线程可以通过 hptr 读写此段内存；
  3. 通过 CUDA API 建立映射 hptr ，返回地址 dptr，设备可以通过 dptr 访问该地址；
  4. 记录信息保存在 struct shmRecvResources* resources->hostHandle 中，资源信息 resources
// 接收设置
NCCLCHECK(transportComm->setup(comm, graph, myInfo, peerInfo, connect, connector, channelId, connIndex));
static ncclResult_t shmRecvSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclPeerInfo* myInfo, struct ncclPeerInfo* peerInfo, struct ncclConnect* connectInfo, struct ncclConnector* recv, int channelId, int connIndex) {
	// 为 resources 申请内存空间
	struct shmRecvResources* resources;
	NCCLCHECK(ncclCalloc(&resources, 1));
	recv->transportResources = resources;
	// 连接信息，准备按照特定格式填充 connectInfo
	struct shmConnectInfo* info = (struct shmConnectInfo*)connectInfo;

	char shmPath[PATH_MAX];
	shmPath[0] = '\0';
	int shmSize = sizeof(struct ncclRecvMem);
	// shmLocality = ncclParamShmLocality();
	// 全局变量 NCCL_SHM_LOCALITY 定义
	// #define SHM_RECV_SIDE 2
	if (shmLocality == SHM_RECV_SIDE) {
		for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) 
			shmSize += comm->buffSizes[p];
	}
	info->shmSize = resources->shmSize = shmSize;
	// 通过 /dev/shm 创建共享内存
	// 保存在 resources->hostMem 中
	// 设备侧操作内存的地址保存在 resources->devHostMem 中
	// resources->hostHandle 保存所有信息
	NCCLCHECK(ncclShmOpen(shmPath, resources->shmSize, (void**)&resources->hostMem, (void**)&resources->devHostMem, 1, &resources->hostHandle));
	TRACE(NCCL_SHM,"Opened shmName %s shmSize %d", shmPath, info->shmSize);
	memcpy(info->shmName, shmPath+sizeof("/dev/shm/nccl-")-1, sizeof(info->shmName));

	return ncclSuccess;
}
Happy_Enger
关注
  • 21
    点赞
  • 踩
  • 19
    收藏
    觉得还不错? 一键收藏
  • 打赏
    打赏
  • 6
    评论
专栏目录
nccl:针对集体多GPU通信的优化原语
04-29
NCCL GPU间通信的优化原语。 介绍 NCCL（发音为“镍”）是用于GPU的标准通信例程的独立库，可实现全缩减，全收集，缩减，广播，缩减分散以及任何基于发送/接收的通讯模式。 它已经过优化，可以在使用PCIe，NVLink，NVswitch的平台以及使用InfiniBand Verbs或TCP / IP套接字的网络上实现高带宽。 NCCL支持在单个节点或多个节点中安装的任意数量的GPU，并且可以在单进程或多进程（例如MPI）应用程序中使用。 有关NCCL使用的更多信息，请参阅。 建造 注意：可以从以下下载NCCL的官方版本和经过测试的版本： : 。 如果选择使用正式版本，则可以跳过以下构建步骤。 要构建库： $ cd nccl $ make -j src.build 如果未在默认的/ usr / local / cuda路径中安装CUDA，则可以使用以下命令定义CUDA路
nccl-tests:NCCL测试
05-04
NCCL测试 这些测试同时检查操作的性能和正确性。 建造 要构建测试，只需键入make 。 如果未在/ usr / local / cuda中安装CUDA，则可以指定CUDA_HOME。 同样，如果未在/ usr中安装NCCL，则可以指定NCCL_HOME。 $ make CUDA_HOME=/path/to/cuda NCCL_HOME=/path/to/nccl NCCL测试依赖MPI来处理多个进程，因此可以处理多个节点。 如果要编译具有MPI支持的测试，则需要设置MPI = 1并将MPI_HOME设置为MPI的安装路径。 $ make MPI=1 MPI_HOME=/path/to/mpi CUDA_HOME=/path/to/cuda NCCL_HOME=/path/to/nccl 用法 NCCL测试可以在多个进程，多个线程以及每个线程多个CUDA设备上运行。 进程数由MP
NCCL源码详解4：bootstrapInit()引导网络bootstrap网络连接建立 视频教程
爱串门的小马驹博客
06-04 1300
Nvidia NCCL bootstrapInit()利用已知的根节(rank0)网络地址(从而获得ncclUniqueId中包含的rank0的网络地址)，建立一个环形网络，allgather获取所有rank的信息。
NVIDIA NCCL 源码学习（八）- 数据通信链路transport的建立
KIDGIN7439的专栏
12-12 6799
上节以ringGraph为例介绍了机器间channel的连接过程，现在环里每个rank都知道了从哪个rank接收数据以及将数据发送给哪个rank，本节具体介绍下P2P和rdma NET场景下数据通信链路的建立过程。
NCCL P2P与共享内存SHM的差异
最新发布
存储随笔
06-13 777
**定义与原理**：在NCCL中提到的SHM，是指在同一个主机节点上，通过操作系统提供的共享内存机制，使得不同进程（或GPU）能够访问同一块物理内存。总结来说，P2P和SHM都是为了提高GPU间通信效率，但P2P更侧重于直接、高效的硬件级数据交换，而SHM则更灵活，能够适应更多种类的系统配置，但可能牺牲一部分性能。- **特点**：P2P通信速度快，延迟低，带宽高，特别适合大规模数据交换，是加速深度学习模型训练和高性能计算中的数据同步（如梯度交换）的关键技术。### 底层原理的差异。
NCCL相关笔记
eternal963的博客
05-18 8646
NCCL是NVIDIA集合通信库（NVIDIA Collective Communications Library）的简称，是用于加速多GPU之间通信的库，能够实现集合通信和点对点通信。
NCCL （NVIDIA Collective Communications Library）
greatcoder的博客
07-25 6237
当n-1次操作完成后，ring-allreduce的第一大步scatter-reduce就已经完成了，此时，第i块gpu的第(i+1)%n块数据已经收集到了所有n块gpu的第(i+1)%n块数据，那么，再进行一次allgather就可以完成算法了。如果卡数比较少的情况下，其实影响不大，但是如果并行的卡很多的时候，就涉及到计算快的卡需要去等待计算慢的卡的情况，造成计算资源的浪费。他也允许点到点的收发通信，包括散播（scatter），聚集（gather），或者all-to-all操作。...
NVIDIA之NCCL：NCCL的简介、安装、使用方法之详细攻略
热门推荐
近期请国内外头部出版社可尽快私信博主！——心比天高，仗剑走天涯，保持热爱，奔赴向梦想！低调，谦虚，自律，反思，成长，还算是比较正能量的博主，公益免费传播……内心特别想在AI界做出一些可以推进历史进程影响力的东西(兴趣使然，有点小情怀，也有点使命感呀)…
06-02 2万+
​ NVIDIA之NCCL：NCCL的简介、安装、使用方法之详细攻略 目录 NCCL的简介 NCCL的安装 NCCL的使用方法 NCCL的案例应用 NCCL的简介 NCCL（NVIDIA Collective Communications Library）是由 NVIDIA 开发的一种高性能的多 GPU 通信库，用于在多个 NVIDIA GPU 之间实现快速的数据传输和协同计算。它可以在深度学习和高性能计算领域中提供分布式训练和数据并行加速的支持。 NCCL实现了针对NVI
【分布式】NCCL部署与测试 - 01
参谋谋的博客
10-07 5408
NCCL简述，NCCL编译测试，代码结构讲解
NCCL基本介绍
qq_43219379的博客
03-14 1万+
NCCL、GPU Direct
nccl-feedstock：nccl的conda-smithy存储库
02-17
关于nccl 主页： ： 软件包许可证：BSD-3-Clause 原料许可证： 简介：用于集体多GPU通信的优化原语 开发： ： 文档： : NVIDIA集体通信库（NCCL）实现了针对NVIDIA GPU性能优化的多GPU和多节点集体通信原语...
mpi-nccl-tests:使用GPU Direct RDMA进行MPI + NCCL测试
05-17
标题中的“mpi-nccl-tests”指的是一个针对MPI（Message Passing Interface）和NCCL（NVIDIA Collective Communications Library）的测试框架，它利用了GPU Direct RDMA（Remote Direct Memory Access）技术来提高...
ubuntu1604+nccl1.2.3+cuda8.0源码
01-04
**标题与描述解析** 标题"ubuntu1604+nccl1.2.3+cuda8.0源码"指的是这个压缩包包含了NVIDIA通信库NCCL（NVIDIA Collective Communications Library）的版本1.2.3的源代码，它是专为在Ubuntu 16.04操作系统上运行...
理解NCCL源码必看：详细论述NCCL源码做了哪些事
这个地瓜真是甜啊
03-21 2823
详解NCCL 源码的实现的功能以及如何实现的
NCCL源码解析: proxy 线程
这个地瓜真是甜啊
02-05 1819
NCCL 源码 Proxy 线程详解
NCCL源码解析④：建图过程
OneFlow深度学习框架
06-27 3141
上次分析了NCCL对机器PCI系统进行拓扑分析的过程，产出的结果为xml格式，接下来，NCCL会根据这个xml进图的建立过程以便之后进行路径搜索。ncclTopoGetSystem的最后会执行ncclTopoGetSystemFromXml将xml格式转成图格式。ncclResult_t ncclTopoGetSystemFromXml(struct ncc...
NVIDIA NCCL 源码学习（十一）- ring allreduce
KIDGIN7439的专栏
10-20 5714
之前的章节里我们看到了nccl send/recv通信的过程，本节我们以ring allreduce为例看下集合通信的过程。整体执行流程和send/recv很像，所以对于相似的流程只做简单介绍，主要介绍ring allreduce自己特有内容。
【分布式】入门级NCCL多机并行实践 - 02
参谋谋的博客
10-20 5156
大模型和分布式训练对数据的吞吐量以及并行度都有很高的要求，NCCL就是在这个背景下诞生的。如果你是一个只会写写Python，调用PyTorch和Horovod的算法萌新，可能对于分布式底层的东西不太了解，在下岗热潮中被主管逼着转变成算子或者通讯库的搬砖工，就会像我一样两眼蒙蔽。因此本文只对自己踩到的坑做一个整理，如果有说错的地方，那就是我说错了。以PyTorch为例，其中spmd接口下的相关定义是用于处理分布式的。但主要是处理单机多CPU情况，因此我们今天只考虑多机（多节点）情况。其中用于实现多节点分布式的
NVIDIA NCCL 源码学习（九）- 单机内ncclSend和ncclRecv的过程
KIDGIN7439的专栏
12-27 1万+
上节介绍了通信链路的建立过程，本节介绍下单机内部ncclSend和ncclRecv的运行过程。 单机内的通信都是通过kernel来完成的，所以整个通信的过程可以分为两步，第一步是准备kernel相关的参数，第二步是实际执行kernel的过程。
NVIDIA NCCL 源码学习
07-28
NVIDIA NCCL（NVIDIA Collective Communications Library）是一种用于高性能并行计算的库，特别适用于多GPU系统中的并行通信操作。如果你想学习NCCL的源码，我可以给你一些建议： 1. 先了解基本概念：在开始研究NCCL源码之前，确保你对并行计算和通信操作有基本的了解。理解NCCL的设计目标和背后的原理会有助于你更好地理解源码。 2. 寻找源码：NCCL的源码可以在NVIDIA的开源GitHub存储库中找到。你可以在https://github.com/NVIDIA/nccl 上找到最新的代码。 3. 阅读文档：NVIDIA提供了NCCL的官方文档，其中包含了详细的API文档和使用指南。在阅读源码之前，先浏览一遍文档，了解库的功能和使用方式，这将有助于你更好地理解源码中的细节。 4. 逐步阅读源码：开始时，可以选择从简单的功能开始阅读，逐步深入到更复杂的部分。从整体架构入手，了解主要的数据结构和函数调用关系。然后，选择一个具体的功能或算法，深入研究相关的源代码。 5. 调试和实践：通过在实际应用中使用NCCL库，你可以更好地理解源码。尝试使用NCCL库进行一些简单的通信操作，并通过调试器进行源码跟踪，观察库的行为和内部工作原理。 6. 参考资料和社区支持：除了官方文档和源码，你还可以参考一些相关的学术论文、博客文章和社区讨论，这些资源可以帮助你更好地理解NCCL的设计和实现。 希望这些建议对你学习NCCL源码有所帮助！如果你有其他问题，请随时提问。
评论 6
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

打赏作者

Happy_Enger

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

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

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

打赏作者

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

抵扣说明：

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

余额充值