NCCL源码解析: proxy 线程

已于 2024-02-06 09:59:32 修改
阅读量1.1k 收藏 9
点赞数 9
分类专栏： NCCL 文章标签： NCCL Linux nvidia
于 2024-02-05 18:16:49 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/u014443578/article/details/136045913
版权

文章目录

前言

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

如有问题，请留言指正。

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

概括

每个GPU对应一个管理线程或者进程，在卡与卡之间建立通信的时候，会额外创建一个代理线程去完成这件事，代理线程是被动的，该做什么事还是由GPU对应的管理线程去通过TCP下发。
代理线程的主要工作有:

  1. 监听TCP端口
  2. 调用 ncclTransportComm 的 proxySharedInit， proxySetup，proxyConnect
  3. 关闭TCP链接

详解

1. 用到的变量

主要关注 comm->proxyState 的初始化，后面会作为理线程参数代使用，用到的时候再来看也行。

// 初始化
commAlloc()
	NCCLCHECK(ncclCalloc(&sharedRes, 1));
bootstrapInit()
	// proxy is aborted through a message; don't set abortFlag
	// 申请内存
	NCCLCHECK(ncclCalloc(&proxySocket, 1));
	// 建立 socket -> proxySocket
	NCCLCHECK(ncclSocketInit(proxySocket, &bootstrapNetIfAddr, comm->magic, ncclSocketTypeProxy, comm->abortFlag));
	// Listen 状态
	NCCLCHECK(ncclSocketListen(proxySocket));
	// 获取地址保存在 state->peerProxyAddresses + rank , IP + Port
	NCCLCHECK(ncclSocketGetAddr(proxySocket, state->peerProxyAddresses+rank));
		struct bootstrapState* state;
		comm->bootstrap = state;
	// 所有节点聚合， state->peerProxyAddresses 保存全部地址
	NCCLCHECK(bootstrapAllGather(state, state->peerProxyAddresses, sizeof(union ncclSocketAddress)));
	// 申请内存初始化 comm->proxyState
	NCCLCHECK(ncclProxyInit(comm, proxySocket, state->peerProxyAddresses));
		NCCLCHECK(ncclCalloc(&comm->sharedRes->proxyState, 1));
		comm->proxyState = comm->sharedRes->proxyState;
		comm->proxyState->refCount = 1;
		comm->proxyState->listenSock = proxySocket;
		comm->proxyState->peerAddresses = state->peerProxyAddresses;

2. proxy 线程创建

主要通过 ncclProxyCreate() 进行 proxyState 对象属性初始化，NCCL 初始化时会创建线程ncclProxyService。

ncclProxyCreate(comm) 
{
	// proxyState 来自 comm->proxyState
	struct ncclProxyState* proxyState = comm->proxyState;
	// 属性初始化，每个属性什么用，用到的时候介绍
	proxyState->tpRank = comm->rank;
	proxyState->tpnRanks = comm->nRanks;
	proxyState->tpLocalnRanks = comm->localRanks;
	proxyState->cudaDev = comm->cudaDev;
	proxyState->abortFlag = comm->abortFlag;
	proxyState->p2pnChannels = comm->p2pnChannels;
	proxyState->p2pChunkSize = comm->p2pChunkSize;
	proxyState->nChannels = comm->nChannels;
	proxyState->allocP2pNetLLBuffers = comm->allocP2pNetLLBuffers;
	proxyState->dmaBufSupport = comm->dmaBufSupport;
	proxyState->ncclNet = comm->ncclNet;
	proxyState->ncclCollNet = comm->ncclCollNet;
	memcpy(proxyState->buffSizes, comm->buffSizes, sizeof(comm->buffSizes));
	// 创建线程
	pthread_create(&comm->proxyState->thread, NULL, ncclProxyService, comm->proxyState);
}
2.1 ncclProxyService()

proxy 服务线程代码, 一个设备起一个 proxy 线程，线程名为 NCCL Service %rank。
线程主要做三件事：

  1. 建立TCP连接
  2. 根据每个卡的线程客户端命令 type 做事
  3. 关闭TCP连接

type 定义如下：

enum ncclProxyMsgType {
  ncclProxyMsgInit = 1,      // 建立 tcp 连接
  ncclProxyMsgSharedInit = 2, // 代理线程调用 ncclTransportComm 的 proxySharedInit
  ncclProxyMsgSetup = 3,   // 代理线程调用 ncclTransportComm 的 proxySetup
  ncclProxyMsgConnect = 4, // 代理线程调用 ncclTransportComm 的 proxyConnect
  ncclProxyMsgStart = 5,   // 还没用
  ncclProxyMsgClose = 6,   // 关闭 TCP 链接
  ncclProxyMsgAbort = 7,   // 还没用
  ncclProxyMsgStop = 8,   // 停用链接，如果所有链接都停用了，代理线程才退出
  ncclProxyMsgConvertFd = 9, // cuMem API support (UDS)
};

线程中主要的处理函数是 proxyServiceInitOp()

// 参数
args = comm->proxyState
void* ncclProxyService(void* _args) {
	struct ncclProxyState* proxyState =  (struct ncclProxyState*) _args;

	// Prepare poll descriptor
	struct ncclProxyConnectionPool connectionPool;
	connectionPool.pools = NULL;
	connectionPool.banks = 0;
	connectionPool.offset = NCCL_PROXY_CONN_POOL_SIZE;

	struct pollfd pollfds[NCCL_MAX_LOCAL_RANKS+1];
	struct ncclProxyLocalPeer peers[NCCL_MAX_LOCAL_RANKS];
	memset(&peers, 0, sizeof(struct ncclProxyLocalPeer)*NCCL_MAX_LOCAL_RANKS);
	for (int s=0; s<NCCL_MAX_LOCAL_RANKS; s++) {
		pollfds[s].fd = -1;
		pollfds[s].events = POLLHUP|POLLIN;
	}
	if (ncclSocketGetFd(proxyState->listenSock, &pollfds[NCCL_MAX_LOCAL_RANKS].fd) != ncclSuccess) {
		WARN("[Proxy Service] Get listenSock fd fails");
		return NULL;
	};
	// 监听输入
	pollfds[NCCL_MAX_LOCAL_RANKS].events = POLLIN;

	int maxnpeers = 0;
	int npeers = 0;
	int stop = 0;
	int asyncOpCount = 0;
	while (stop == 0 || (stop == 1 && npeers > 0)) {
		/* Even if local comm aborts, we cannot let proxy thread exit if we still have peer
			* connections. Need to wait until all other related comms call abort and safely exit
			* together, or we could face segmentation fault. */
		// 本地退出，不能推出线程，需要等其他 comms 也停止才能一起退出
		if (*proxyState->abortFlag != 0) stop = 1;
		/* never let proxy service thread blocks in poll, or it cannot receive abortFlag. */
		int ret;
		do {
			ret = poll(pollfds, NCCL_MAX_LOCAL_RANKS+1, asyncOpCount ? 0 : 500);
		} while (ret < 0 && errno == EINTR);
		if (ret < 0) {
			WARN("[Proxy Service] Poll failed: %s", strerror(errno));
			return NULL;
		}
		if (pollfds[NCCL_MAX_LOCAL_RANKS].revents) {
			int s = 0;
			while (s < NCCL_MAX_LOCAL_RANKS && pollfds[s].fd >= 0) s++;
			if (s == NCCL_MAX_LOCAL_RANKS) {
				WARN("[Proxy service] Too many connections (%d max)", NCCL_MAX_LOCAL_RANKS);
				return NULL;
			}
			if (maxnpeers < s+1) maxnpeers = s+1;
			// 初始化 socket
			if (ncclSocketInit(&peers[s].sock) != ncclSuccess) {
				WARN("[Service thread] Initialize peers[%d].sock fails", s);
				return NULL;
			}
			// accept
			if (ncclSocketAccept(&peers[s].sock, proxyState->listenSock) != ncclSuccess) {
				WARN("[Service thread] Accept failed %s", strerror(errno));
			} else {
				// 监听 fd 到 pollfds
				if (ncclSocketGetFd(&peers[s].sock, &pollfds[s].fd) != ncclSuccess) {
					WARN("[Service thread] Get peers[%d].sock fd fails", s);
					return NULL;
				}
				npeers++;
				peers[s].tpLocalRank = -1;
			}
		}
		for (int s=0; s<maxnpeers; s++) {
			struct ncclProxyLocalPeer* peer = peers+s;
			struct ncclSocket* sock = &peer->sock;
			int closeConn = 0;
			int type = 0;
			ncclResult_t res = ncclSuccess;
			if (pollfds[s].fd == -1) 
				continue;

			// Progress all ops for this ncclProxyLocalPeer
			ncclProxyAsyncOp* op = peer->asyncOps;
			while (op != nullptr) {
				ncclProxyAsyncOp* opnext = op->next; /* in case op is freed in proxyProgressAsync */
				type = op->type;
				res = proxyProgressAsync(op, proxyState, &asyncOpCount, peer, &connectionPool);
				if (res == ncclSuccess || res == ncclInProgress) {
					op = opnext;
				} else {
					// Res is a bad result
					closeConn = 1;
					WARN("[Service thread] Error encountered progressing operation=%s, res=%d, closing connection", ncclProxyMsgTypeStr[type], res);
					break;
				}
			}

			// Check for additional ops coming in
			// 检查是否有输入
			if (pollfds[s].revents & POLLIN) {
				int closed;
				// 先接收 Type
				res = ncclSocketTryRecv(sock, &type, sizeof(int), &closed, false /*blocking*/);
				if (res != ncclSuccess && res != ncclInProgress) {
					WARN("[Service thread] Could not receive type from localRank %d, res=%u, closed=%d", peer->tpLocalRank, res, closed);
					closeConn = 1;
				} else if (closed) {
					INFO(NCCL_INIT|NCCL_NET|NCCL_PROXY, "[Service thread] Connection closed by localRank %d", peer->tpLocalRank);
					closeConn = 1;
				} else if (res == ncclSuccess) { // We received something from the sock
					// 接收到数据，根据 type 做不同的动作
					if (type == ncclProxyMsgStop) {
					// 关闭连接
						stop = 1;
					closeConn = 1;
					} else if (type == ncclProxyMsgClose) {
					// 关闭连接
						closeConn = 1;
					} else if (proxyMatchOpType(type)) {
						// 处理客户端即设备的请求，根据 type 进行不同的处理
						res = proxyServiceInitOp(type, peers+s, &connectionPool, proxyState, &asyncOpCount);
					} else {
						// 关闭连接
						WARN("[Service thread] Unknown command %d from localRank %d", type, peer->tpLocalRank);
						closeConn = 1;
					}

					INFO(NCCL_PROXY, "Received and initiated operation=%s res=%d", ncclProxyMsgTypeStr[type], res);
				}
			} else if (pollfds[s].revents & POLLHUP) {
				// 关闭连接
				closeConn = 1;
			}

			if (res != ncclSuccess && res != ncclInProgress) {
				// 关闭连接
				WARN("[Proxy Service %d] Failed to execute operation %s from rank %d, retcode %d", proxyState->tpRank, ncclProxyMsgTypeStr[type], peer->tpRank, res);
				closeConn = 1;
			}

			if (closeConn) {
				// 关闭连接
				ncclSocketClose(sock);

				if (op != nullptr) {
					asyncProxyOpDequeue(peer, op);
					asyncOpCount--;
				}
				pollfds[s].fd = -1;
				npeers--;
			}
		}
	}

	// 退出操作
	// Wait for all operations to complete and stop progress thread before freeing any resource
	if (ncclProxyProgressDestroy(proxyState) != ncclSuccess) {
		WARN("[Proxy Service] proxyDestroy failed");
	}
	for (int s=0; s<maxnpeers; s++) {
		ncclSocketClose(&peers[s].sock);
	}
	ncclProxyFreeConnections(&connectionPool, proxyState);
	ncclSocketClose(proxyState->listenSock);
	free(proxyState->listenSock);
	proxyOpsFree(proxyState);
	return NULL;
}
2.2 proxyServiceInitOp()

线程中的主要处理函数，因为客户端发送数据的时候是先发什么后发什么的顺序，所以接收也先按一定的顺序接收数据，然后调用 proxyProgressAsync 进行处理；


// 本地 rank 的 proxyState
// peers 是保存在服务端的数据，数据保存的是客户端的信息
// peer 抽象的是客户端对象
res = proxyServiceInitOp(type, peers+s, &connectionPool, proxyState, &asyncOpCount);
static ncclResult_t proxyServiceInitOp(int type, struct ncclProxyLocalPeer* peer, struct ncclProxyConnectionPool* connectionPool, struct ncclProxyState* proxyState, int* asyncOpCount) {
	// 服务端 sock
	struct ncclSocket* sock = &peer->sock;
	// 申请内存
	struct ncclProxyAsyncOp* asyncOp;
	NCCLCHECK(ncclCalloc(&asyncOp, 1));

	asyncOp->type = type;
	// 按照客户端发送的顺序，接收各个字段
	// 接收 connection， 指向发送端 connection 对象的首地址
	NCCLCHECK(ncclSocketRecv(sock, &asyncOp->connection, sizeof(void*)));
	// 获取发送长度
	NCCLCHECK(ncclSocketRecv(sock, &asyncOp->reqSize, sizeof(int)));
	// 获取接收缓冲区大小
	NCCLCHECK(ncclSocketRecv(sock, &asyncOp->respSize, sizeof(int)));
	if (asyncOp->reqSize) {
		// 如果发送长度大于0，发送端会发送数据，接收端要接收数据
		// 先申请内存再接收数据
		NCCLCHECK(ncclCalloc(&asyncOp->reqBuff, asyncOp->reqSize));
		NCCLCHECK(ncclSocketRecv(sock, asyncOp->reqBuff, asyncOp->reqSize));
	}

	// Store opId for completion response
	// 接收发送端 opId 的首地址
	NCCLCHECK(ncclSocketRecv(sock, &asyncOp->opId, sizeof(asyncOp->opId)));

	// 如果发送端要接收数据，则接收数据大小大于0，服务端要申请发送缓冲区内存
	if (asyncOp->respSize) 
		NCCLCHECK(ncclCalloc(&asyncOp->respBuff, asyncOp->respSize));

	// 请求 asyncOp 加入peer 对象链表中 peer->asyncOps
	asyncProxyOpEnqueue(peer, asyncOp);

	(*asyncOpCount)++;
	// 处理请求
	NCCLCHECK(proxyProgressAsync(asyncOp, proxyState, asyncOpCount, peer, connectionPool));
	return ncclSuccess;
}
2.2 proxyProgressAsync()

处理请求函数，根据参数 type 进行不同的逻辑处理，然后按照一定的顺序返回数据

NCCLCHECK(proxyProgressAsync(asyncOp, proxyState, asyncOpCount, peer, connectionPool));
static ncclResult_t proxyProgressAsync(struct ncclProxyAsyncOp* op, struct ncclProxyState* proxyState, int* asyncOpCount, struct ncclProxyLocalPeer* peer, struct ncclProxyConnectionPool* connectionPool) {
	int done = 1;
	if (op->type == ncclProxyMsgSetup) {
		// 调用 proxy proxySetup API
		TRACE(NCCL_PROXY, "proxyProgressAsync::proxySetup() opId=%p", op->opId);
		NCCLCHECK(op->connection->tcomm->proxySetup(op->connection, proxyState, op->reqBuff, op->reqSize, op->respBuff, op->respSize, &done));
	} else if (op->type == ncclProxyMsgConnect) {
		// 调用 proxy proxyConnect API
		TRACE(NCCL_PROXY, "proxyProgressAsync::proxyConnect() opId=%p op.reqBuff=%p", op->opId, op->reqBuff);
		NCCLCHECK(op->connection->tcomm->proxyConnect(op->connection, proxyState, op->reqBuff, op->reqSize, op->respBuff, op->respSize, &done));
	} else if (op->type == ncclProxyMsgSharedInit) {
		int nChannels = (int) *op->reqBuff;
		// 调用 proxy proxySharedInit API
		TRACE(NCCL_PROXY, "proxyProgressAsync::ncclProxyMsgSharedInit opId=%p op.reqBuff=%p nChannels=%d", op->opId, op->reqBuff, nChannels);
		if (op->connection->tcomm->proxySharedInit) NCCLCHECK(op->connection->tcomm->proxySharedInit(op->connection, proxyState, nChannels));
		__atomic_store_n(&op->connection->state, connSharedInitialized, __ATOMIC_RELEASE);
	} else if (op->type == ncclProxyMsgConvertFd) {
		int fd = *(int *)op->reqBuff;
		TRACE(NCCL_PROXY, "proxyProgressAsync::ncclProxyMsgConvertFd opId=%p op.reqBuff=%p fd=%d", op->opId, op->reqBuff, fd);
		NCCLCHECK(proxyConvertFd(peer, op->opId, proxyState, fd)); // cuMem API support
	} else if (op->type == ncclProxyMsgInit) {
		// 
		TRACE(NCCL_PROXY, "proxyProgressAsync::ncclProxyMsgInit opId=%p op.reqBuff=%p", op->opId, op->reqBuff);
		NCCLCHECK(proxyConnInit(peer, connectionPool, proxyState, (ncclProxyInitReq*) op->reqBuff, (ncclProxyInitResp*) op->respBuff, &op->connection));
		static ncclResult_t proxyConnInit(struct ncclProxyLocalPeer* peer, struct ncclProxyConnectionPool* connectionPool, struct ncclProxyState* proxyState, ncclProxyInitReq* req, ncclProxyInitResp* resp, struct ncclProxyConnection** connection) 
		{
			int id;
			// 为 connectionPool-> pools 分配空间，
			// connectionPool->offset++
			// id = ((pool->banks-1) << NCCL_PROXY_CONN_POOL_SIZE_POW2) + pool->offset;
			// offset 为 (1 << 7) 个，为一个 bank
			NCCLCHECK(ncclProxyNewConnection(connectionPool, &id));
			// 根据 id 获取 bank 与 offset
			// 根据 bank与 offset 获取 ncclProxyConnection 首地址 connection
			NCCLCHECK(ncclProxyGetConnection(connectionPool, id, connection));
			// 填充 connection
			(*connection)->sock = &peer->sock;
			(*connection)->transport = req->transport;
			(*connection)->send = req->send;
			(*connection)->tpLocalRank = req->tpLocalRank;
			(*connection)->sameProcess = req->sameProcess;
			peer->tpLocalRank = req->tpLocalRank;
			peer->tpRank = req->tpRank;
			// connection 首地址给 resp->connection, 要告诉客户端
			resp->connection = *connection;

			(*connection)->tcomm = (*connection)->send ? &ncclTransports[(*connection)->transport]->send : &ncclTransports[(*connection)->transport]->recv;
			// If we need proxy progress, let's allocate ops and start the thread
			if ((*connection)->tcomm->proxyProgress) {
				NCCLCHECK(proxyProgressInit(proxyState));
				struct ncclProxyProgressState* state = &proxyState->progressState;
				strncpy(resp->devShmPath, state->opsPoolShmSuffix, sizeof(resp->devShmPath));
			}
			INFO(NCCL_NET|NCCL_PROXY, "New proxy %s connection %d from local rank %d, transport %d", (*connection)->send ? "send":"recv", id, (*connection)->tpLocalRank, (*connection)->transport);
			__atomic_store_n(&(*connection)->state, connInitialized, __ATOMIC_RELEASE);
			return ncclSuccess;
		}
	} else 
		return ncclInternalError;

	if (done) {
		INFO(NCCL_PROXY, "proxyProgressAsync opId=%p op.type=%d op.reqBuff=%p op.respSize=%d done", op->opId, op->type, op->reqBuff, op->respSize);
		if (op->type == ncclProxyMsgSetup)
			__atomic_store_n(&op->connection->state, connSetupDone, __ATOMIC_RELEASE);
		else if (op->type == ncclProxyMsgConnect)
			__atomic_store_n(&op->connection->state, connConnected, __ATOMIC_RELEASE);
		/* if setup or connect is done, we should not return any error at this point since
			* ncclSocketSend might already send the respBuff to the requester. If we still choose
			* to abort and close the connection, it can cause segfault if the requester is using
			* the respBuff. */

		// Send the opId for referencing async operation
		// 发送 opId
		NCCLCHECK(ncclSocketSend(op->connection->sock, &op->opId, sizeof(op->opId)));

		// Send the response size
		// 发送接收大小
		NCCLCHECK(ncclSocketSend(op->connection->sock, &op->respSize, sizeof(op->respSize)));

		if (op->respSize) {
			// Send the response
			// 发送响应
			NCCLCHECK(ncclSocketSend(op->connection->sock, op->respBuff, op->respSize));
		}
		// op 移出链表
		asyncProxyOpDequeue(peer, op);
		(*asyncOpCount)--;
		return ncclSuccess;

	} else if (*proxyState->abortFlag != 0) {
		return ncclInternalError;
	}

	return ncclInProgress;
}

4. ncclProxyConnect()

以其中链接为例 ：如果要使用代理，那么首先要先连接，通过 type 为 ncclProxyMsgInit 告诉代理，我要链接，代理线程会 accept 建立 socket, 返回连接的 ncclProxyConnection connection 对象的首地址

链接流程如下，主要关注数据传输，有的传数据，有的传首地址：


// p2p send connector
// rank GPU 设备连接 proxy TCP 服务端，服务端建立保存连接，申请通信所需的内存资源
struct ncclConnector* send
NCCLCHECK(ncclProxyConnect(comm, TRANSPORT_P2P, 1, tpProxyRank, &send->proxyConn));
ncclResult_t ncclProxyConnect(struct ncclComm* comm, int transport, int send, int tpProxyRank, struct ncclProxyConnector* proxyConn) {
  struct ncclSocket* sock;
  int ready, proxyRank = -1;
  struct ncclProxyState* sharedProxyState = comm->proxyState;

  // Keep one connection per mlocal rank
  for (int i = 0; i < comm->localRanks; ++i) {
    /* find the proxy rank in comm. */
    if (comm->topParentRanks[comm->localRankToRank[i]] == tpProxyRank) {
      proxyRank = comm->localRankToRank[i];
      break;
    }
  }
  proxyConn->sameProcess = comm->peerInfo[proxyRank].pidHash == comm->peerInfo[comm->rank].pidHash ? 1 : 0;
  // Keep one connection per local rank
  proxyConn->connection = NULL;
  proxyConn->tpRank = tpProxyRank;
  // peerSocks 初始化
  if (sharedProxyState->peerSocks == NULL) {
    NCCLCHECK(ncclCalloc(&sharedProxyState->peerSocks, comm->sharedRes->tpNLocalRanks));
    NCCLCHECK(ncclCalloc(&sharedProxyState->proxyOps, comm->sharedRes->tpNLocalRanks));
    NCCLCHECK(ncclCalloc(&sharedProxyState->sharedDevMems, comm->sharedRes->tpNLocalRanks));
    for (int i = 0; i < comm->sharedRes->tpNLocalRanks; ++i) {
      NCCLCHECK(ncclSocketSetFd(-1, &sharedProxyState->peerSocks[i]));
    }
  }

  proxyConn->tpLocalRank = comm->sharedRes->tpRankToLocalRank[proxyConn->tpRank];
  sock = sharedProxyState->peerSocks + proxyConn->tpLocalRank;
  NCCLCHECK(ncclSocketReady(sock, &ready));
  if (!ready) {
	// scoket 初始化 socket
    NCCLCHECK(ncclSocketInit(sock, sharedProxyState->peerAddresses+proxyConn->tpRank, comm->sharedRes->magic, ncclSocketTypeProxy, comm->abortFlag));
	// 连接 proxy 服务线程中监听的端口
    NCCLCHECK(ncclSocketConnect(sock));
  }

  struct ncclProxyInitReq req = {0};
  req.transport = transport;
  req.send = send;
  req.tpLocalRank = comm->topParentLocalRanks[comm->localRank];
  req.tpRank = comm->topParentRanks[comm->rank];
  req.sameProcess = proxyConn->sameProcess;

  struct ncclProxyInitResp resp = {0};
  // This usually sends proxyConn->connection to identify which connection this is.
  // However, this is part of the response and therefore is ignored
  // 收发消息初始化，proxy 服务端申请内存，建立连接
  NCCLCHECK(ncclProxyCallBlocking(comm, proxyConn, ncclProxyMsgInit, &req, sizeof(req), &resp, sizeof(resp)));
  // resp.connection 为服务端的 connection 对象的首地址
  proxyConn->connection = resp.connection;

  // If we need proxy progress, map progress ops
  struct ncclTransportComm* tcomm = send ? &ncclTransports[transport]->send : &ncclTransports[transport]->recv;
  if (tcomm->proxyProgress) {
    char poolPath[] = "/dev/shm/nccl-XXXXXX";
    strncpy(poolPath+sizeof("/dev/shm/nccl-")-1, resp.devShmPath, sizeof("XXXXXX")-1);
    struct ncclProxyOps* proxyOps = sharedProxyState->proxyOps + proxyConn->tpLocalRank;
    if (proxyOps->pool == NULL) {
      NCCLCHECK(ncclShmOpen(poolPath, sizeof(struct ncclProxyOpsPool), (void**)(&proxyOps->pool), NULL, 0, &proxyOps->handle));
      proxyOps->nextOps = proxyOps->nextOpsEnd = proxyOps->freeOp = -1;
    }
  }
  INFO(NCCL_NET|NCCL_PROXY, "Connection to proxy localRank %d -> connection %p", proxyConn->tpLocalRank, proxyConn->connection);
  return ncclSuccess;
}
4.1 ncclProxyCallBlocking()

调用代理线程接口，即开始发送命令，接收返回。

// 客户端通知 proxy 服务端调用响应接口，服务端根据 type 做不同的处理
// ncclProxyMsgInit 表示服务端进行通信初始化
NCCLCHECK(ncclProxyCallBlocking(comm, proxyConn, ncclProxyMsgInit, &req, sizeof(req), &resp, sizeof(resp)));
ncclResult_t ncclProxyCallBlocking(struct ncclComm* comm, struct ncclProxyConnector* proxyConn, int type, void* reqBuff, int reqSize, void* respBuff, int respSize) {
	// Alloc some memory to act as a handle
	ncclResult_t res = ncclSuccess;
	void* opId = malloc(1);
	// ncclProxyCallAsync()
	// 首先发送 type
	// 再发送 proxyConn->connection 的首地址
	// 发送 reqSize
	// 发送 respSize
	// 如果 reqSize 大于0，说明有发送数据，即发送数据
	// 发送 opId 的首地址
	NCCLCHECKGOTO(ncclProxyCallAsync(comm, proxyConn, type, reqBuff, reqSize, respSize, opId), res, fail);
		struct ncclProxyState* sharedProxyState = comm->proxyState;
		sock = sharedProxyState->peerSocks + proxyConn->tpLocalRank;
		// 将当前 请求放入 state 的链表中  state->expectedResponses;
		NCCLCHECK(expectedProxyResponseEnqueue(sharedProxyState, opId, respSize));
		{
			struct ncclExpectedProxyResponse* ex;
			NCCLCHECK(ncclCalloc(&ex, 1));
			ex->opId = opId;

			// Pre-alloc response buffer
			ex->respBuff = malloc(respSize);
			ex->respSize = respSize;
			ex->done     = false;
			struct ncclExpectedProxyResponse* list = state->expectedResponses;
			if (list == NULL) {
				state->expectedResponses = ex;
				return ncclSuccess;
			}
			while (list->next) list = list->next;
			list->next = ex;
		}
		    

	do {
		res = ncclPollProxyResponse(comm, proxyConn, respBuff, opId);
		{
			int found = 0;
			// 如果 opId 在链表中找到，且 done 字段已被置为 True, 则拷贝数据到 respBuff， found 置 1
  			NCCLCHECK(expectedProxyResponseDequeue(sharedProxyState, opId, respBuff, &found));
		}
	} while (res == ncclInProgress);

exit:
	free(opId);
	return res;
fail:
	goto exit;
}
4.2 ncclPollProxyResponse()

发送的时候有 opId 作为此次通信的标识，代理线程返回数据时也会把这个opId带回来
所以接受的时候要比较 opId, 如果与本次发送的 opId 一样，那么就接收成功；
如果不一样，那么把接受的数据放入缓冲区，继续接收

// 轮询等待 opId 的返回数据
res = ncclPollProxyResponse(comm, proxyConn, respBuff, opId);
ncclResult_t ncclPollProxyResponse(struct ncclComm* comm, struct ncclProxyConnector* proxyConn, void* respBuff, void* opId) {
	struct ncclProxyState* sharedProxyState = comm->proxyState;
	// Receive the connection pointer from the Proxy
	// 检查停止字段
	if (*comm->abortFlag) {
		WARN("Comm %p is in abort state", comm);
		return ncclInternalError;
	}

	if (sharedProxyState->peerSocks == NULL) 
		return ncclInternalError;

	// Check response queue
	int found = 0;
	// 如果 opId 在链表中找到，且 done 字段已被置为 True, 则拷贝数据到 respBuff， found 置 1
	NCCLCHECK(expectedProxyResponseDequeue(sharedProxyState, opId, respBuff, &found));
	if (found == 0) {
		// 发送完之后，还没收到回复，虽然有 opId, 但是 done 字段仍为 False, 所以 found == 0
		// Attempt to read in a new response header from the proxy thread
		// 对于没有父节点的 comm来说，tpLocalRank 就是 comm->localrank
		// 获取发送端的 socket
		struct ncclSocket* sock = sharedProxyState->peerSocks + proxyConn->tpLocalRank;

		void* recvOpId;
		int offset = 0;
		// 接收数据，先接受 opId
		if (ncclSuccess != ncclSocketProgress(NCCL_SOCKET_RECV, sock, &recvOpId, sizeof(recvOpId), &offset)) {
			WARN("Socket recv failed while polling for opId=%p", opId);
			return ncclInternalError;
		}

		// 确保接收全部数据， offset == 0 返回 ncclInProgress 继续接收数据
		if (offset == 0) {
			return ncclInProgress;
		// If we've returned a partial response, block to receive the rest of it
		} else if (offset < sizeof(recvOpId)) {
			while (offset < sizeof(recvOpId))
			NCCLCHECK(ncclSocketProgress(NCCL_SOCKET_RECV, sock, &recvOpId, sizeof(recvOpId), &offset));
		}

		INFO(NCCL_PROXY, "ncclPollProxyResponse Received new opId=%p", recvOpId);

		// Now do a blocking recv of the response size
		int respSize = 0;
		// 接收返回数据的大小
		NCCLCHECK(ncclSocketRecv(sock, &respSize, sizeof(respSize)));

		// If there's a respSize to recv
		if (respSize > 0) {
			// 有返回数据
			if (recvOpId != opId) {
				// Unexpected response, need to buffer the socket data
				// 对于意想不到的 opId, 申请内存保存数据
				respBuff = malloc(respSize);
			}
			assert(respBuff != NULL);
			// 接收返回的数据
			NCCLCHECK(ncclSocketRecv(sock, respBuff, respSize));
		}

		if (recvOpId == opId) {
			// 如果已经接收了 opId 的数据，则在 state->expectedResponses 链表中移除 opId 相对应的项
			INFO(NCCL_PROXY, "recvOpId=%p matches expected opId=%p", recvOpId, opId);
			NCCLCHECK(expectedProxyResponseRemove(sharedProxyState, recvOpId));
			// 返回成功
			return ncclSuccess;
		} else {
			INFO(NCCL_PROXY, "Queuing opId=%p respBuff=%p respSize=%d", recvOpId, respBuff, respSize);
			// Store the result and mark response as completed
			// 如果接收的是其他 opId 的数据，则拷贝数据到缓冲区，并置 elem->done 为 True
			NCCLCHECK(expectedProxyResponseStore(sharedProxyState, recvOpId, respBuff, respSize));
			// 返回，继续处理接收数据
			return ncclInProgress;
		}
	} else {
		INFO(NCCL_PROXY, "ncclPollProxyResponse Dequeued cached opId=%p", opId);
	}

	return ncclSuccess;
}
Happy_Enger
关注
  • 9
    点赞
  • 踩
  • 9
    收藏
    觉得还不错? 一键收藏
  • 打赏
    打赏
  • 3
    评论
专栏目录
mpi-nccl-tests:使用GPU Direct RDMA进行MPI + NCCL测试
05-17
MPI和NCCL GPU直接RDMA测试 建造 mkdir build && cd build cmake .. && make
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路
NVIDIA NCCL 源码学习（十三）- IB SHARP
KIDGIN7439的专栏
03-20 1706
因此基于这一点，mellanox提出了SHARP，将计算offload到了IB switch，每个节点只需要发送一次数据，这块数据会被交换机完成规约，然后每个节点再接收一次就得到了完整结果。
理解NCCL源码必看：详细论述NCCL源码做了哪些事
最新发布
这个地瓜真是甜啊
03-21 1235
详解NCCL 源码的实现的功能以及如何实现的
NCCL源码解析: 共享内存连接的建立
这个地瓜真是甜啊
02-06 1193
NCCL 源码中两个设备通过共享内存通信初始化
NVIDIA NCCL 源码学习（十二）- double binary tree
KIDGIN7439的专栏
12-22 2808
上节我们以ring allreduce为例看到了集合通信的过程，但是随着训练任务中使用的gpu个数的扩展，ring allreduce的延迟会线性增长，为了解决这个问题，NCCL引入了tree算法，即double binary tree。
nccl 源码分析 从 ncclAllReduce 的执行开始认识nccl源代码
eloudy的专栏
01-04 802
这个info是在ncclAllReduce()中定义的 struct ncclInfo info，其中info->comm 是ncclAllReduce 传递进来的第五个参数 ncclComm* comm。所以，启动相关的阿cuda kernel等，应该是在后面的 ncclGroupEnd() 中，通过解析 comm->tasks的数据元素来启动的。那么需要看一下 ncclIntruQueueEnqueue 到底对t中的 t->op做了什么解析，t->op是这个函数的第二个参数的op成员；
NVIDIA NCCL 源码学习（十）- 多机间ncclSend和ncclRecv的过程
KIDGIN7439的专栏
06-01 5341
通信由kernel和proxy线程协调完成，send端kernel负责将数据从input搬运到buf，proxy线程负责将buf中数据通过网络发送给recv端，kernel和proxy间通过队列实现生产者消费者模式，send端通过rdma send发送数据，和recv端通过队列实现生产者消费者模式，队列位于send端，recv端每次下发一个wr到rq之后会执行rdma write通知send端
NVIDIA NCCL 源码学习（九）- 单机内ncclSend和ncclRecv的过程
KIDGIN7439的专栏
12-27 8497
上节介绍了通信链路的建立过程，本节介绍下单机内部ncclSend和ncclRecv的运行过程。 单机内的通信都是通过kernel来完成的，所以整个通信的过程可以分为两步，第一步是准备kernel相关的参数，第二步是实际执行kernel的过程。
NVIDIA NCCL 源码学习（十一）- ring allreduce
KIDGIN7439的专栏
10-20 4120
之前的章节里我们看到了nccl send/recv通信的过程，本节我们以ring allreduce为例看下集合通信的过程。整体执行流程和send/recv很像，所以对于相似的流程只做简单介绍，主要介绍ring allreduce自己特有内容。
nccl-master.rar
04-09
深度学习依赖库nccl
nccl_2.4.8-1+cuda10.0_x86_64.txz
09-17
当使用paddle多GPU时报错，缺少NCCL，将文件解压后cp include/nccl.h /home/myname/cuda/include/ cp /lib/libnccl* /home/myname/cuda/lib64/ 即可。
nccl-repo-ubuntu1804-2.4.2-ga-cuda10.1_1-1_amd64.deb
06-09
nccl
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
ubuntu1604+nccl1.2.3+cuda8.0源码
01-04
NVIDIA多GPU加速工具nccl的源码，适用于Ubuntu16.04+cuda8.0环境，安装方式参照readme文件
nccl-feedstock：nccl的conda-smithy存储库
02-17
关于nccl 主页： ： 软件包许可证：BSD-3-Clause 原料许可证： 简介：用于集体多GPU通信的优化原语 开发： ： 文档： : NVIDIA集体通信库（NCCL）实现了针对NVIDIA GPU性能优化的多GPU和多节点集体通信原语...
NVIDIA NCC​L 源码学习（三）- 机器内拓扑分析
KIDGIN7439的专栏
09-28 5326
上节介绍到所有节点执行了bootstrap网络连接的建立，接下来介绍下拓扑分析。由于gpu机器架构是多种多样的，一台机器上可能有多个网卡，多个gpu卡，卡间连接也各不相同，因此需要对机器内设备连接拓扑进行分析以达到各种拓扑结构下性能都尽可能的好。
NVIDIA NCCL 源码学习（八）- 数据通信链路transport的建立
KIDGIN7439的专栏
12-12 5570
上节以ringGraph为例介绍了机器间channel的连接过程，现在环里每个rank都知道了从哪个rank接收数据以及将数据发送给哪个rank，本节具体介绍下P2P和rdma NET场景下数据通信链路的建立过程。
torch.distributed.DistBackendError: NCCL error in: ../torch/csrc/distributed/c10d/ProcessGroupNCCL.cpp:1275, internal error, NCCL version 2.14.3 ncclInternalError: Internal check failed. Last error: Proxy Call to rank 0 failed (Connect)
05-27
这是一个 PyTorch 分布式训练过程中遇到的错误。出现这个错误的原因可能是因为在使用 NCCL 通信时，进程间的连接失败了。解决这个问题的方法有： 1. 确认你的 NCCL 版本是否正确，建议使用与 PyTorch 版本匹配的 NCCL 版本。 2. 确认你的网络环境是否正常，特别是防火墙是否允许进程间通信。 3. 确认你的代码是否正确，特别是分布式训练时是否正确设置了进程间通信的参数。 如果以上方法都无法解决问题，建议查看 PyTorch 和 NCCL 的文档，或者在 GitHub 上搜索相关的 issue。

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 3
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

打赏作者

Happy_Enger

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

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

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

打赏作者

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

抵扣说明：

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

余额充值