A test for the core assumption of data + signal of MSCCL++

[saeedmaleki]

A test that can check if our core assumption about data->write followed by signal()/wait() always preserve the write order on the receiving GPU side.

Something like this:

__device__ void pingPongTest(volatile int* sendBuff, mscclpp::channel::SimpleDeviceChannel devChan, int rank, int worldSize, int nranksPerNode,
                           int remoteRank, size_t nelemsPerGPU) {
  int nTries = 1000;
  int flusher = 0;
  int rank1Offset = 10000000;
  if (true){
    for (int i = 0; i < nTries; i++){
      if (rank == 0){
        if (i > 0){
          if (threadIdx.x == 0)
            devChan.wait();
          __syncthreads();
          for (int j = threadIdx.x; j < nelemsPerGPU; j+=blockDim.x){
            if (sendBuff[j] != rank1Offset+i-1+j){
              printf("rank 0 ERROR: sendBuff[%d] = %d, expected %d\n", j, sendBuff[j], 100000+i-1+j);
            }
          }
        }
        for (int j = threadIdx.x; j < nelemsPerGPU; j+=blockDim.x){
          sendBuff[j] = i+j;
        }
        __syncthreads();
        // __threadfence_system(); // not necessary if we make sendBuff volatile
        if (threadIdx.x == 0)
          devChan.putWithSignal(0, nelemsPerGPU * sizeof(int));
      }
      if (rank == 1){
        if (threadIdx.x == 0)
          devChan.wait();
        __syncthreads();
        for (int j = threadIdx.x; j < nelemsPerGPU; j+=blockDim.x){
          if (sendBuff[j] != i+j){
            printf("rank 1 ERROR: sendBuff[%d] = %d, expected %d\n", j, sendBuff[j], i+j);
          }
        }
        if (i < nTries-1){
          for (int j = threadIdx.x; j < nelemsPerGPU; j+=blockDim.x){
            sendBuff[j] = rank1Offset+i+j;
          }
          __syncthreads();
          // __threadfence_system(); // not necessary if we make sendBuff volatile
          if (threadIdx.x == 0)
            devChan.putWithSignal(0, nelemsPerGPU * sizeof(int));
        }
      }
      flusher++;
      if (flusher == 100){
        devChan.flush();
        flusher = 0;
      }
    }
  }
}

[chhwang]

Hi @saeedmaleki, I'm trying to add this to our unit test, and the code is here. Seems this is an orthogonal issue, but the test code returns an error message NonblockingFuture::get() called before ready (Mscclpp failure: InvalidUsage). Do you have any ideas why? The commands is:

mpirun --allow-run-as-root -tag-output -np 2 ./build/test/mp_unit_tests --gtest_filter=ChannelOneToOneTest.PingPongIb

[chhwang]

The issue is resolved and the test is added to unit tests in #81