[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)) {
        // 双方都需要注册 注册后可以直接访问的对端的内存地址就是regAddr
        NCCLCHECK(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;