[4.2-1] NCCL新版本的register如何实现的？

0. overall

会包括四节内容，先跳过了上层的调用部分，这部分的内容会在第五章节更新。简言之就是上层的torch调用empty等接口创建tensor，新版的话会通过ncclMemAlloc和ncclCommRegister对这个tensor的地址开辟和注册。当tensor在tp/pp之间传输的时候就会直接使用注册好的内存地址实现zerocopy，完成数据的高带宽传输。
这里的4.2-1到4.2-2两节会从enqueue.cc内怎么实现zerocopy的角度去看原生两个gpu之间如何通过注册完成zerocopy行为的。

1. addP2pToplan()

大致是send recv传输操作加到执行计划内：

1. 判断当前是机内还是机间
2. 注册buffer，机间就是 ncclRegisterP2pNetBuffer，机内就是 ncclRegisterP2pIpcBuffer，这里会去transport层找自己的register buffer的具体实现
3. 生成设备工作结构，ncclDevWorkP2p
4. 为非自发送创建proxy操作
   机间注册buffer的时候会多一步判断，必须得保证PSN是关闭的。并且会给每个channel都注册buffer。

bool pxnUsed = !ncclPxnDisable(comm) && comm->isAllNvlink && comm->maxLocalRanks > 1;
if (bytes[dir] > 0 && proxySameProcess[dir] && protocol[dir] == NCCL_PROTO_SIMPLE && (!pxnUsed)) {int regFlag = 0;NCCLCHECK(ncclCalloc(&handles[dir], nChannelsMax));for (int part = 0; part < nChannelsMax; part++) {int channelId = ncclP2pChannelForPart(comm->p2pnChannels, base, part);struct ncclChannelPeer** channelPeers = comm->channels[channelId].peers;int peerRank = dir ? sendRank : recvRank;struct ncclConnector* conn = dir ? &channelPeers[peerRank]->send[connIndex]: &channelPeers[peerRank]->recv[connIndex];if (conn->conn.flags & NCCL_DIRECT_NIC)ncclRegisterP2pNetBuffer(comm, addrs[dir], bytes[dir], conn, &regFlag, &handles[dir][part], &plan->cleanupQueue);if (!regFlag) break;}netRegistered[dir] = regFlag ? true : false;}

机内的话发现，有实际数据，有地址是simple协议和非selfsend就会去注册ipcbuffer。注册具体流程：ncclRegisterP2pIpcBuffer在[4.2-2]内有详细说明（因为这一块需要单开一个说明不同进程不同gpu怎么拿到对方注册地址的）内（fd和handle去交换拿到的注册buffer）：
注册后拿到对端注册的地址 regAddr ，这里举例说比如机内rank0发给rank1。那么rank0是sender，对应这里时dir=1的情况，sender把自己的地址addr[1]放进ncclRegisterP2pIpcBuffer 内拿到对端注册的 regAddr ，但是它 sendAddr = regAddr，也就是说把对端注册的buffer地址覆盖了自己的sendAddr。同理receiver也是这样，通过完成注册后打印看到刚刚的举例真实情况如下：
Rank 0 (发送方dir=1):

• sendBytes=0x400000 (4MB数据)
• sendRank=1, sendAddr=0x7fa275000000 (本地发送缓冲区地址)
• recvRank=1, recvAddr=(nil) (接收地址为空，因为是发送方，这不是0，这是p2pSchedule初始化成1的)
• ipcRegistered[0]=0, ipcRegistered[1]=1 (只有发送方向注册了IPC)
  Rank 1 (接收方dir=0):
• sendBytes=0xffffffffffffffff (无发送数据)
• sendRank=0, sendAddr=(nil) (发送地址为空，因为是接收方)
• recvRank=0, recvAddr=0x40ca00000 (本地接收缓冲区地址)
• ipcRegistered[0]=1, ipcRegistered[1]=0 (只有接收方向注册了IPC)

else if (bytes[dir] > 0 && addrs[dir] && protocol[dir] == NCCL_PROTO_SIMPLE && !selfSend) {int peerRank = dir ? sendRank : recvRank;int regFlag = 0;int channelId = ncclP2pChannelForPart(comm->p2pnChannels, base, 0);struct ncclChannelPeer** channelPeers = comm->channels[channelId].peers;struct ncclConnector* conn = dir ? &channelPeers[peerRank]->send[connIndex]: &channelPeers[peerRank]->recv[connIndex];void* regAddr = NULL;// [NCCL_P2P_WRITE] 表示可以写入对端内存// [NCCL_P2P_READ]  表示可以从对端内存读取if (conn->conn.flags & (NCCL_P2P_WRITE | NCCL_P2P_READ)) {// 双方都需要注册 注册后可以直接访问的对端的内存地址就是regAddrNCCLCHECK(ncclRegisterP2pIpcBuffer(comm, addrs[dir], bytes[dir], peerRank, &regFlag, &regAddr, &plan->cleanupQueue));if (regFlag) {if (dir == 0 && (conn->conn.flags & NCCL_P2P_WRITE)) recvAddr = regAddr;else if (dir == 1 && (conn->conn.flags & NCCL_P2P_READ)) sendAddr = regAddr;}}ipcRegistered[dir] = regFlag ? true : false;}

selfsend就是单进程，只下一个proxyOp。非selfsend就是涉及进程间的传输，就下两个proxyop，dir=0是recv，dir=1是send。

struct ncclProxyOp proxyOps[2] = {};
int nProxyOps = selfSend ? 0 : 2;