add doc string

Author: Binyang2014

include/mscclpp/proxy_channel.hpp

@@ -207,11 +207,24 @@ struct ProxyChannel {
     put(dst, offset, src, offset, size);
   }
 
+  /// @brief Push a @ref TriggerData to the FIFO.
+  /// @param dst The destination memory region.
+  /// @param dstOffset The offset into the destination memory region.
+  /// @param src The source memory region.
+  /// @param srcOffset The offset into the source memory region.
+  /// @param width The width of the 2D region.
+  /// @param height The height of the 2D region.
   __forceinline__ __device__ void put2D(MemoryId dst, uint64_t dstOffset, MemoryId src, uint64_t srcOffset,
                                         uint32_t width, uint32_t height) {
     fifo_.push(ChannelTrigger(TriggerData, dst, dstOffset, src, srcOffset, width, height, semaphoreId_).value);
   }
 
+  /// @brief Push a @ref TriggerData to the FIFO.
+  /// @param dst The destination memory region.
+  /// @param src The source memory region.
+  /// @param offset The common offset into the destination and source memory regions.
+  /// @param width The width of the 2D region.
+  /// @param height The height of the 2D region.
   __forceinline__ __device__ void put2D(MemoryId dst, MemoryId src, uint64_t offset, uint32_t width, uint32_t height) {
     put2D(dst, offset, src, offset, width, height);
   }
@@ -232,6 +245,13 @@ struct ProxyChannel {
     fifo_.push(ChannelTrigger(TriggerData | TriggerFlag, dst, dstOffset, src, srcOffset, size, semaphoreId_).value);
   }
 
+  /// Push a @ref TriggerData and a @ref TriggerFlag at the same time to the FIFO.
+  /// @param dst The destination memory region.
+  /// @param dstOffset The offset into the destination memory region.
+  /// @param src The source memory region.
+  /// @param srcOffset The offset into the source memory region.
+  /// @param width The width of the 2D region.
+  /// @param height The height of the 2D region.
   __forceinline__ __device__ void put2DWithSignal(MemoryId dst, uint64_t dstOffset, MemoryId src, uint64_t srcOffset,
                                                   uint32_t width, uint32_t height) {
     fifo_.push(
@@ -247,6 +267,12 @@ struct ProxyChannel {
     putWithSignal(dst, offset, src, offset, size);
   }
 
+  /// Push a @ref TriggerData and a @ref TriggerFlag at the same time to the FIFO.
+  /// @param dst The destination memory region.
+  /// @param src The source memory region.
+  /// @param offset The common offset into the destination and source memory regions.
+  /// @param width The width of the 2D region.
+  /// @param height The height of the 2D region.
   __forceinline__ __device__ void put2DWithSignal(MemoryId dst, MemoryId src, uint64_t offset, uint32_t width,
                                                   uint32_t height) {
     put2DWithSignal(dst, offset, src, offset, width, height);
@@ -328,6 +354,11 @@ struct SimpleProxyChannel {
     proxyChan_.put(dst_, dstOffset, src_, srcOffset, size);
   }
 
+  /// Push a @ref TriggerData to the FIFO.
+  /// @param dstOffset The offset into the destination memory region.
+  /// @param srcOffset The offset into the source memory region.
+  /// @param width The width of the 2D region.
+  /// @param height The height of the 2D region.
   __forceinline__ __device__ void put2D(uint64_t dstOffset, uint64_t srcOffset, uint32_t width, uint32_t height) {
     proxyChan_.put2D(dst_, dstOffset, src_, srcOffset, width, height);
   }
@@ -348,6 +379,11 @@ struct SimpleProxyChannel {
     proxyChan_.putWithSignal(dst_, dstOffset, src_, srcOffset, size);
   }
 
+  /// Push a @ref TriggerData and a @ref TriggerFlag at the same time to the FIFO.
+  /// @param dstOffset The offset into the destination memory region.
+  /// @param srcOffset The offset into the source memory region.
+  /// @param width The width of the 2D region.
+  /// @param height The height of the 2D region.
   __forceinline__ __device__ void put2DWithSignal(uint64_t dstOffset, uint64_t srcOffset, uint32_t width,
                                                   uint32_t height) {
     proxyChan_.put2DWithSignal(dst_, dstOffset, src_, srcOffset, width, height);
@@ -358,6 +394,10 @@ struct SimpleProxyChannel {
   /// @param size The size of the transfer.
   __forceinline__ __device__ void putWithSignal(uint64_t offset, uint64_t size) { putWithSignal(offset, offset, size); }
 
+  /// Push a @ref TriggerData, a @ref TriggerFlag, and a @ref TriggerSync at the same time to the FIFO.
+  /// @param offset The common offset into the destination and source memory regions.
+  /// @param width The width of the 2D region.
+  /// @param height The height of the 2D region.
   __forceinline__ __device__ void put2DWithSignal(uint64_t offset, uint32_t width, uint32_t height) {
     put2DWithSignal(offset, offset, width, height);
   }

src/connection.cc

@@ -66,12 +66,10 @@ void CudaIpcConnection::write2D(RegisteredMemory dst, uint64_t dstOffset, Regist
   char* dstPtr = (char*)dst.data();
   char* srcPtr = (char*)src.data();
 
-  INFO(MSCCLPP_P2P,
-       "CudaIpcConnection write: from %p to %p, width %lu height %lu dstPitch %lu srcPitch %lu dstOffset %lu srcOffset "
-       "%lu",
-       srcPtr + srcOffset, dstPtr + dstOffset, width, height, dst.pitch(), src.pitch(), dstOffset, srcOffset);
   MSCCLPP_CUDATHROW(cudaMemcpy2DAsync(dstPtr + dstOffset, dst.pitch(), srcPtr + srcOffset, src.pitch(), width, height,
                                       cudaMemcpyDeviceToDevice, stream_));
+  INFO(MSCCLPP_P2P, "CudaIpcConnection write: from %p to %p, width %lu height %lu dstPitch %lu srcPitch %lu",
+       srcPtr + srcOffset, dstPtr + dstOffset, width, height, dst.pitch(), src.pitch());
 }
 
 void CudaIpcConnection::updateAndSync(RegisteredMemory dst, uint64_t dstOffset, uint64_t* src, uint64_t newValue) {

test/mp_unit/proxy_channel_tests.cu

@@ -69,10 +69,10 @@ void ProxyChannelOneToOneTest::setupMeshConnections(
 
 __constant__ DeviceHandle<mscclpp::SimpleProxyChannel> gChannelOneToOneTestConstProxyChans;
 
-__device__ size_t getTileElementOffset(int elementId, int width, int rowIndex, int colIndex, int nElemInPitch) {
+__device__ size_t getTileElementOffset(int elementId, int width, int rowIndex, int colIndex, int nElemPerPitch) {
   int rowIndexInTile = elementId / width;
   int colIndexInTile = elementId % width;
-  return (rowIndex + rowIndexInTile) * nElemInPitch + (colIndex + colIndexInTile);
+  return (rowIndex + rowIndexInTile) * nElemPerPitch + (colIndex + colIndexInTile);
 }
 
 __global__ void kernelProxyTilePingPong(int* buff, int rank, int pitch, int rowIndex, int colIndex, int width,
@@ -83,14 +83,14 @@ __global__ void kernelProxyTilePingPong(int* buff, int rank, int pitch, int rowI
   int flusher = 0;
   size_t offset = rowIndex * pitch + colIndex * sizeof(int);
   size_t nElem = width * hight;
-  size_t nElemPerInPitch = pitch / sizeof(int);
+  size_t nElemPerPitch = pitch / sizeof(int);
   for (int i = 0; i < nTries; i++) {
     if (rank == 0) {
       if (i > 0) {
         if (threadIdx.x == 0) proxyChan.wait();
         __syncthreads();
         for (int j = threadIdx.x; j < nElem; j += blockDim.x) {
-          size_t tileOffset = getTileElementOffset(j, width, rowIndex, colIndex, nElemPerInPitch);
+          size_t tileOffset = getTileElementOffset(j, width, rowIndex, colIndex, nElemPerPitch);
           if (sendBuff[tileOffset] != offset + i - 1 + j) {
             // printf("rank 0 ERROR: sendBuff[%d] = %d, expected %d\n", j, sendBuff[j], rank1Offset + i - 1 + j);
             *ret = 1;
@@ -99,7 +99,7 @@ __global__ void kernelProxyTilePingPong(int* buff, int rank, int pitch, int rowI
         }
       }
       for (int j = threadIdx.x; j < nElem; j += blockDim.x) {
-        size_t tileOffset = getTileElementOffset(j, width, rowIndex, colIndex, nElemPerInPitch);
+        size_t tileOffset = getTileElementOffset(j, width, rowIndex, colIndex, nElemPerPitch);
         sendBuff[tileOffset] = i + j;
       }
       __syncthreads();
@@ -110,7 +110,7 @@ __global__ void kernelProxyTilePingPong(int* buff, int rank, int pitch, int rowI
       if (threadIdx.x == 0) proxyChan.wait();
       __syncthreads();
       for (int j = threadIdx.x; j < nElem; j += blockDim.x) {
-        size_t tileOffset = getTileElementOffset(j, width, rowIndex, colIndex, nElemPerInPitch);
+        size_t tileOffset = getTileElementOffset(j, width, rowIndex, colIndex, nElemPerPitch);
         if (sendBuff[tileOffset] != i + j) {
           // printf("rank 1 ERROR: sendBuff[%d] = %d, expected %d\n", j, sendBuff[j], i + j);
           *ret = 1;
@@ -119,7 +119,7 @@ __global__ void kernelProxyTilePingPong(int* buff, int rank, int pitch, int rowI
       }
       if (i < nTries - 1) {
         for (int j = threadIdx.x; j < nElem; j += blockDim.x) {
-          size_t tileOffset = getTileElementOffset(j, width, rowIndex, colIndex, nElemPerInPitch);
+          size_t tileOffset = getTileElementOffset(j, width, rowIndex, colIndex, nElemPerPitch);
           sendBuff[tileOffset] = offset + i + j;
         }
         __syncthreads();