Add tile-based API

[Binyang2014]

Provide tile-based api

void put2D(uint64_t dstOffset, uint64_t srcOffset, uint32_t width, uint32_t height)
void put2DWithSignal(uint64_t dstOffset, uint64_t srcOffset, uint32_t width, uint32_t height)

To support this, add a new structure fields2D in ChannelTrigger. In this structure we replace the 64bit size to two 32 bit fields (a 32bit width and a 32bit height). Also add another flag multiDimensionFlag in fields2D to distinguish with fields structure

Example to use tile-based API:
When setup the connections, need to call channelService->addPitch first

  for (int r = 0; r < worldSize; r++) {
    if (r == rank) {
      continue;
    }
    std::shared_ptr<mscclpp::Connection> conn;
    if ((rankToNode(r) == rankToNode(gEnv->rank)) && !useIbOnly) {
      conn = communicator->connectOnSetup(r, 0, mscclpp::Transport::CudaIpc);
    } else {
      conn = communicator->connectOnSetup(r, 0, ibTransport);
    }
    connections[r] = conn;
    communicator->sendMemoryOnSetup(recvBufRegMem, r, 0);
    auto remoteMemory = communicator->recvMemoryOnSetup(r, 0);
    communicator->setup();

    mscclpp::SemaphoreId cid = channelService->add2DChannel(conn, std::pair<size_t, size_t>(dstPitch, srcPitch));
    communicator->setup();
  }

Then use the put2D API

if (threadIdx.x == 0) proxyChan.put2DWithSignal(offset, width * sizeof(int), height);

[Binyang2014]

The API cudaMemcpy2DAsync seems slower than cudaMemcpyAsync for 1D data. Need to investigate

[saeedmaleki]

This doesn't make much sense to me, why do we need an extra an way of building a semaphore? We only need to provide a 2D write over 1D arrays. So, just a 2D write is enough. Right?

[Binyang2014]

Since we need to set pitch/stride for the channel (the name semaphore is not accurate). The reason we don't set stride in the put2D API is our trigger is only 128bit. We don't have extra bits for the it.