NCCL源码解析: P2P 连接的建立

2024-02-07 13:20
文章标签 源码 连接 解析 建立 p2p nccl

本文主要是介绍NCCL源码解析: P2P 连接的建立，希望对大家解决编程问题提供一定的参考价值，需要的开发者们随着小编来一起学习吧！

文章目录

  • 前言
  • 概括
  • 详解
    • ncclTransportP2pSetup()

前言

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
};

本文重点关注 P2P。

接口如下：

struct ncclTransport p2pTransport = {"P2P",p2pCanConnect,{ p2pSendSetup, p2pSendConnect, p2pSendFree, NULL, p2pSendProxySetup, NULL, p2pSendProxyFree, NULL },{ p2pRecvSetup, p2pRecvConnect, p2pRecvFree, NULL, p2pRecvProxySetup, NULL, p2pRecvProxyFree, NULL }
};

发送建立流程为 p2pCanConnect() -> p2pSendSetup() -> p2pSendProxySetup()
接收建立流程为 p2pCanConnect() -> p2pRecvSetup() -> p2pRecvProxySetup()

先检查两个卡支不支持 P2P，主要检查两项：设备支不支持、路径支不支持，路径类型要小于 PATH_PXB，即不通过主桥的路径；
然后 p2pSendSetup() 填充一下 p2pConnectInfo, 向 proxy 线程请求 ncclProxyMsgSetup;
proxy 线程调用 p2pSendProxySetup(), 在本卡内申请显存，返回首地址以及相应的句柄devIpc, 其他进程或者线程可以通过这个句柄，获得此显存的操作地址。（我猜是让对端 GPU 卡也来操作这段内存，从而完成通信，还没看到那一步，完事来更新）
接收与发送机制一样。

详解

ncclTransportP2pSetup()

此P2P 非彼 P2P。ncclTransportP2pSetup 的 P2P 是广义上的两个设备之间的通信设置，包含 P2P、网络以及共享内存等。
建立两卡通信的入口函数。
因为要与 proxy 双线程操作，但是又是同步的，所以下文把两个线程的操作线性展开了，请注意。

ncclTransportP2pSetup(comm, &ringGraph, 0)
ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, int connIndex, int* highestTransportType/*=NULL*/)
{// 信息保存在 data[i] 中， i 为 rank// data[i] 大小为 2 * 64 个 connect, 先存放 recv, 再存放 sendrecvData[i] = data[i];// recvData[]  所有 recvChannels 的 ncclConnect 缓冲区的首地址// 首先 <0> 表示处理的是接收，处理与前一个 rank 的连接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(transportComm->setup(comm, graph, myInfo, peerInfo, connect, connector, channelId, connIndex)){NCCLCHECK(ncclCalloc(&resources, 1));recv->transportResources = resources;struct p2pConnectInfo* info = (struct p2pConnectInfo*)connectInfo;// 如果使用nvlink, 且两个GPU 计算能力一样，(gpu1->gpu.cudaCompCap == 80), 那么 useRead = 1// 如果通过参数 P2P_READ_ENABLE 设置该值，  P2P 使用 read 而不是 writeinfo->read = useRead; for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) if (!(info->read && p == NCCL_PROTO_SIMPLE)) recvSize += comm->buffSizes[p];// 如果同一个进程内的，且 DirectDisable 没有设置，P2P_USE_CUDA_MEMCPY 参数没有设置，并且ncclCuMemEnable 为假// 那么{resources->type = P2P_DIRECT;recv->conn.flags |= info->read ? NCCL_DIRECT_READ : NCCL_DIRECT_WRITE;}// 建立到 proxy 的连接, 连接信息在 recv->proxyConnNCCLCHECK(ncclProxyConnect(comm, TRANSPORT_P2P, 0, tpProxyRank, &recv->proxyConn));// 请求 proxy 执行 ncclProxyMsgSetup// recvSize += comm->buffSizes[p];// 发送数据 4字节 recvSize  10485760 = 4096 + NCCL_NUM_PROTOCOLS 3 类型的缓冲区大小// 接收数据缓冲区 info->p2pBuff// 要接收的大小 sizeof(struct ncclP2pBuff)// info->p2pBuff 保存 buf 信息NCCLCHECK(ncclProxyCallBlocking(comm, &recv->proxyConn, ncclProxyMsgSetup, &recvSize, sizeof(int), &info->p2pBuff, sizeof(struct ncclP2pBuff)));// 下面为 proxy 线程// proxy 线程接收数据进行处理{// op->connection : 设备与 proxy 连接的控制对象// proxyState : rank 的 ncclProxyState// op->reqBuff : proxy 本地的接收缓冲区首地址， 按照  op->reqSize 大小申请// op->reqSize ： 客户端发送的发送数据的大小// op->respBuff: proxy 本地的发送缓冲区的首地址，按照 op->respSize 大小申请// p2pRecvProxySetup ： 设备申请内存，首地址信息存入 respBuffNCCLCHECK(op->connection->tcomm->proxySetup(op->connection, proxyState, op->reqBuff, op->reqSize, op->respBuff, op->respSize, &done));static ncclResult_t p2pRecvProxySetup(struct ncclProxyConnection* connection, struct ncclProxyState* proxyState, void* reqBuff, int reqSize, void* respBuff, int respSize, int* done) {// 获取设备侧告知的 recvSize 的值 10485760int size = *((int*)reqBuff);struct ncclP2pBuff* p2pBuff = (struct ncclP2pBuff*)respBuff;NCCLCHECK(ncclP2pAllocateShareableBuffer(size, &p2pBuff->ipcDesc, &p2pBuff->directPtr));ncclResult_t ncclP2pAllocateShareableBuffer(size_t size, ncclIpcDesc *ipcDesc, void **ptr) {// 在设备侧申请内存，地址保存在 ptrNCCLCHECK(ncclCudaCalloc((char **)ptr, size));// cudaIpcGetMemHandle ： 获取现有设备内存分配的进程间内存句柄// 获取指向使用cudaMalloc创建的现有设备内存分配的基址的指针，并将其导出以供另一个进程使用// __host__ cudaError_t cudaIpcGetMemHandle ( cudaIpcMemHandle_t* handle, void* devPtr )// 获取现有设备内存分配的进程间内存句柄。// 参数:// handle - 指向用户分配的 cudaIpcMemHandle 以返回句柄的指针。// devPtr - 指向先前分配的设备内存的基指针cudaError_t res = cudaIpcGetMemHandle(&ipcDesc->devIpc, *ptr);// cudaIpcOpenMemHandle : 打开从另一个进程导出的进程间内存句柄并返回可用于本地进程的设备指针// __host__ cudaError_t cudaIpcOpenMemHandle ( void** devPtr, cudaIpcMemHandle_t handle, unsigned int  flags )
​// 打开从另一个进程导出的进程间内存句柄并返回可用于本地进程的设备指针。// 参数// devPtr - 返回设备指针// handle - cudaIpcMemHandle 打开// flags - 此操作的标志。必须指定为cudaIpcMemLazyEnablePeerAccess}p2pBuff->size = size;connection->transportResources = p2pBuff->directPtr;}}// 下面不是 proxy 线程// 设备收到 proxy 返回的信息： 设备内部申请的缓冲区首地址，以及地址句柄 ipcDesc->devIpc// comm->peerInfo AllGather1 时保存的所有 rank 的信息: rank cudaDev hostHash pidHash busId// info->rank = myInfo->rank// p2pBuff : info->p2pBuff// devMem : (void**)&resources->recvDevMem 设备接收资源的接收缓冲区内存指针地址// ipcPtr : &resources->recvMemIpc 设备接收资源的接收 内存Ipc 指针地址NCCLCHECK(p2pMap(comm, myInfo, comm->peerInfo+info->rank, &info->p2pBuff, (void**)&resources->recvDevMem, &resources->recvMemIpc));static ncclResult_t p2pMap(struct ncclComm *comm, struct ncclPeerInfo* myInfo, struct ncclPeerInfo* peerInfo, struct ncclP2pBuff* p2pBuff, void** devMem, void** ipcPtr){// 如果 ncclCuMemEnable 为假，且两个 GPU 设备在同一进程中// 那么{// 如果本设备与对端设备不是同一设备if (peerInfo->cudaDev != myInfo->cudaDev) {// 如果可以从设备直接访问 peerDevice，则可以通过调用 cudaDeviceEnablePeerAccess() 来启用访问cudaError_t err = cudaDeviceEnablePeerAccess(peerInfo->cudaDev, 0);}// 把 proxy 从设备申请的内存首地址赋值给 *devMem，即 resources->recvDevMem// resources->recvDevMem = p2pBuff->directPtr;*devMem = p2pBuff->directPtr;// 同一个设备不用 ipc*ipcPtr = NULL;}else{if ((myInfo->pidHash == peerInfo->pidHash) && (peerInfo->cudaDev == myInfo->cudaDev)) {// 同一个进程，同一个设备// Same PID and GPU*devMem = p2pBuff->directPtr;*ipcPtr = NULL;} else {// 不同进程或者不同设备// Different PID or different GPUNCCLCHECK(ncclP2pImportShareableBuffer(comm, comm->topParentRanks[peerInfo->rank], p2pBuff->size, &p2pBuff->ipcDesc, devMem));ncclResult_t ncclP2pImportShareableBuffer(struct ncclComm *comm, int tpPeer, size_t size, ncclIpcDesc *ipcDesc, void **devMemPtr) {// cudaIpcOpenMemHandle : 打开从另一个进程导出的进程间内存句柄并返回可用于本地进程的设备指针// __host__ cudaError_t cudaIpcOpenMemHandle ( void** devPtr, cudaIpcMemHandle_t handle, unsigned int  flags )
​// 打开从另一个进程导出的进程间内存句柄并返回可用于本地进程的设备指针。// 参数// devPtr - 返回设备指针// handle - cudaIpcMemHandle 打开// flags - 此操作的标志。必须指定为cudaIpcMemLazyEnablePeerAccess// 通过 ipcDesc->devIpc 获取设备内存首地址 devMemPtrCUDACHECK(cudaIpcOpenMemHandle(devMemPtr, ipcDesc->devIpc, cudaIpcMemLazyEnablePeerAccess));}// devMem 已经赋值为设备内存首地址*ipcPtr = *devMem;}}}}}// 发送// 信息保存在 data[i] 中， i 为 rank// data[i] 大小为 2 * 64 个 connect, 先存放 recv, 再存放 send// sendData[]  所有 recvChannels 的发送 ncclConnect 缓冲区的首地址sendData[i] = recvData[i] + recvChannels;// 调用发送，处理与后一个 rank 的连接NCCLCHECKGOTO(selectTransport<1>(comm, graph, sendData[i]+sendChannels++, c, sendPeer, connIndex, &type), ret, fail);static ncclResult_t selectTransport(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclConnect* connect, int channelId, int peer, int connIndex, int* transportType) {NCCLCHECK(transportComm->setup(comm, graph, myInfo, peerInfo, connect, connector, channelId, connIndex));ncclResult_t p2pSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclPeerInfo* myInfo, struct ncclPeerInfo* peerInfo, struct ncclConnect* connectInfo, struct ncclConnector* send, int channelId, int connIndex){NCCLCHECK(ncclCalloc(&resources, 1));send->transportResources = resources;info->read = useRead;if (graph && connIndex == 1) info->read = 0;const char* useReadStr = info->read ? "/read" : "";// For P2P Read the SIMPLE buffer is tagged on the end of the ncclSendMem structureif (info->read) // 只有读的时候，使用缓冲区 NCCL_PROTO_SIMPLEsendSize += comm->buffSizes[NCCL_PROTO_SIMPLE];info->rank = myInfo->rank;resources->type = P2P_DIRECT;send->conn.flags |= info->read ? NCCL_DIRECT_READ : NCCL_DIRECT_WRITE;// 与接收一样的操作// 设备收到 proxy 返回的信息保存在 p2pBuff中： 设备内部申请的缓冲区首地址，以及地址句柄 ipcDesc->devIpcNCCLCHECK(ncclProxyCallBlocking(comm, &send->proxyConn, ncclProxyMsgSetup, &sendSize, sizeof(int), &info->p2pBuff, sizeof(struct ncclP2pBuff)));// p2pMap : 根据接收到的信息做一个发送缓冲区的首地址解析，得到 sendDevMem 或者 sendMemIpc// info->rank = myInfo->rank// p2pBuff : info->p2pBuff// devMem : (void**)&resources->recvDevMem 设备接收资源的接收缓冲区内存指针地址// ipcPtr : &resources->recvMemIpc 设备接收资源的接收 内存Ipc 指针地址NCCLCHECK(p2pMap(comm, myInfo, comm->peerInfo+info->rank, &info->p2pBuff, (void**)&resources->sendDevMem, &resources->sendMemIpc));}}if (sendPeer == recvPeer) {if (recvChannels+sendChannels) {NCCLCHECKGOTO(bootstrapSend(comm->bootstrap, recvPeer, bootstrapTag, data[i], sizeof(struct ncclConnect)*(recvChannels+sendChannels)), ret, fail);NCCLCHECKGOTO(bootstrapRecv(comm->bootstrap, recvPeer, bootstrapTag, data[i], sizeof(struct ncclConnect)*(recvChannels+sendChannels)), ret, fail);sendData[i] = data[i];recvData[i] = data[i]+sendChannels;}} else {// 如果 sendPeer recvPeer 不是同一个// 假设 0 -> 1 -> 2, 当前 rank 为 1// sendPeer = 2， recvPeer = 0// sendPeer ： 我作为发送的对端 rank// recvPeer ： 我作为接收的对端 rankif (recvChannels) // 向前一个 rank 发送 recvChannels 个接收连接信息NCCLCHECKGOTO(bootstrapSend(comm->bootstrap, recvPeer, bootstrapTag, recvData[i], sizeof(struct ncclConnect)*recvChannels), ret, fail);if (sendChannels) // 向后一个 rank 发送 recvChannels 个发送连接信息NCCLCHECKGOTO(bootstrapSend(comm->bootstrap, sendPeer, bootstrapTag, sendData[i], sizeof(struct ncclConnect)*sendChannels), ret, fail);if (sendChannels) // 接收后一个 rank 的接收连接信息到 sendDataNCCLCHECKGOTO(bootstrapRecv(comm->bootstrap, sendPeer, bootstrapTag, sendData[i], sizeof(struct ncclConnect)*sendChannels), ret, fail);if (recvChannels) // 接收前一个 rank 的接收发送信息NCCLCHECKGOTO(bootstrapRecv(comm->bootstrap, recvPeer, bootstrapTag, recvData[i], sizeof(struct ncclConnect)*recvChannels), ret, fail);}
}

这篇关于NCCL源码解析: P2P 连接的建立的文章就介绍到这儿，希望我们推荐的文章对编程师们有所帮助！


http://www.chinasem.cn/article/687863。

相关文章

网页解析 lxml 库--实战

网页解析 lxml 库--实战

lxml库使用流程 lxml 是 Python 的第三方解析库，完全使用 Python 语言编写，它对 XPath表达式提供了良好的支 持，因此能够了高效地解析 HTML/XML 文档。本节讲解如何通过 lxml 库解析 HTML 文档。 pip install lxml lxm| 库提供了一个 etree 模块，该模块专门用来解析 HTML/XML 文档，下面来介绍一下 lxml 库
阅读更多...
W外链微信推广短连接怎么做？

W外链微信推广短连接怎么做？

制作微信推广链接的难点分析 一、内容创作难度 制作微信推广链接时，首先需要创作有吸引力的内容。这不仅要求内容本身有趣、有价值，还要能够激起人们的分享欲望。对于许多企业和个人来说，尤其是那些缺乏创意和写作能力的人来说，这是制作微信推广链接的一大难点。 二、精准定位难度 微信用户群体庞大，不同用户的需求和兴趣各异。因此，制作推广链接时需要精准定位目标受众，以便更有效地吸引他们点击并分享链接
阅读更多...
JAVA智听未来一站式有声阅读平台听书系统小程序源码

JAVA智听未来一站式有声阅读平台听书系统小程序源码

智听未来，一站式有声阅读平台听书系统 🌟&nbsp;开篇：遇见未来，从“智听”开始 在这个快节奏的时代，你是否渴望在忙碌的间隙，找到一片属于自己的宁静角落？是否梦想着能随时随地，沉浸在知识的海洋，或是故事的奇幻世界里？今天，就让我带你一起探索“智听未来”——这一站式有声阅读平台听书系统，它正悄悄改变着我们的阅读方式，让未来触手可及！ 📚&nbsp;第一站：海量资源，应有尽有 走进“智听
阅读更多...
【C++】_list常用方法解析及模拟实现

【C++】_list常用方法解析及模拟实现

相信自己的力量，只要对自己始终保持信心，尽自己最大努力去完成任何事，就算事情最终结果是失败了，努力了也不留遗憾。💓💓💓 目录   ✨说在前面 🍋知识点一：什么是list？ •🌰1.list的定义 •🌰2.list的基本特性 •🌰3.常用接口介绍 🍋知识点二：list常用接口 •🌰1.默认成员函数 🔥构造函数（⭐） 🔥析构函数 •🌰2.list对象
阅读更多...
Java ArrayList扩容机制 （源码解读）

Java ArrayList扩容机制 （源码解读）

结论：初始长度为10，若所需长度小于1.5倍原长度，则按照1.5倍扩容。若不够用则按照所需长度扩容。 一. 明确类内部重要变量含义         1：数组默认长度         2:这是一个共享的空数组实例，用于明确创建长度为0时的ArrayList ，比如通过 new ArrayList<>(0)，ArrayList 内部的数组 elementData 会指向这个 EMPTY_EL
阅读更多...
如何在Visual Studio中调试.NET源码

如何在Visual Studio中调试.NET源码

今天偶然在看别人代码时，发现在他的代码里使用了Any判断List<T>是否为空。 我一般的做法是先判断是否为null，再判断Count。 看了一下Count的源码如下： 1 [__DynamicallyInvokable]2 public int Count3 {4 [__DynamicallyInvokable]5 get
阅读更多...
工厂ERP管理系统实现源码（JAVA）

工厂ERP管理系统实现源码（JAVA）

工厂进销存管理系统是一个集采购管理、仓库管理、生产管理和销售管理于一体的综合解决方案。该系统旨在帮助企业优化流程、提高效率、降低成本，并实时掌握各环节的运营状况。 在采购管理方面，系统能够处理采购订单、供应商管理和采购入库等流程，确保采购过程的透明和高效。仓库管理方面，实现库存的精准管理，包括入库、出库、盘点等操作，确保库存数据的准确性和实时性。 生产管理模块则涵盖了生产计划制定、物料需求计划、
阅读更多...
Java 连接Sql sever 2008

Java 连接Sql sever 2008

Java 连接Sql sever 2008 /Sql sever 2008 R2 import java.sql.Connection; import java.sql.DriverManager; import java.sql.ResultSet; import java.sql.Statement; public class TestJDBC
阅读更多...
OWASP十大安全漏洞解析

OWASP十大安全漏洞解析

OWASP（开放式Web应用程序安全项目）发布的“十大安全漏洞”列表是Web应用程序安全领域的权威指南，它总结了Web应用程序中最常见、最危险的安全隐患。以下是对OWASP十大安全漏洞的详细解析： 1. 注入漏洞（Injection） 描述：攻击者通过在应用程序的输入数据中插入恶意代码，从而控制应用程序的行为。常见的注入类型包括SQL注入、OS命令注入、LDAP注入等。 影响：可能导致数据泄
阅读更多...
建立升序链表

建立升序链表

题目1181：遍历链表 时间限制：1 秒 内存限制：32 兆 特殊判题：否 提交：2744 解决：1186 题目描述： 建立一个升序链表并遍历输出。 输入： 输入的每个案例中第一行包括1个整数：n(1<=n<=1000)，接下来的一行包括n个整数。 输出： 可能有多组测试数据，对于每组数据， 将n个整数建立升序链表，之后遍历链表并输出。 样例输
阅读更多...

Copyright China编程 23002807@qq.com

沪ICP备2023033072号-2