added stream synch after hipMemset (#609)

[ROCm/rccl commit: 06bce9d0c9]
2022-08-30 16:18:37 -06:00
parent f18868f439
commit 2cd63dac42
5 changed files with 5 additions and 3 deletions
@@ -93,6 +93,7 @@ static ncclResult_t ncclCudaCallocDebug(const char *filefunc, int line, T** ptr,
  CUDACHECK(hipStreamDestroy(stream));
 #else
  CUDACHECK(hipMemset(*ptr, 0, nelem*sizeof(T)));
+  CUDACHECK(hipStreamSynchronize(NULL));
 #endif
  INFO(NCCL_ALLOC, "%s:%d Cuda Alloc Size %ld pointer %p", filefunc, line, nelem*sizeof(T), *ptr);
  int dev;
@@ -99,6 +99,7 @@ namespace RcclUnitTesting
      ERROR("Unable to call hipMemset\n");
      return TEST_FAIL;
    }
+    hipStreamSynchronize(NULL);
    return TEST_SUCCESS;
  }

@@ -172,7 +172,7 @@ int main(int argc, char* argv[])
      HIPCHECK(hipMalloc(recvbuff+i, size*sizeof(float)));
      h_recvbuff[i] = (float*) malloc (size *sizeof(float));
      HIPCHECK(hipMemset(recvbuff[i], 0, size*sizeof(float)));
-
+      HIPCHECK(hipStreamSynchronize(NULL));
      HIPCHECK(hipStreamCreate(s+i));
  }

@@ -174,7 +174,7 @@ int main(int argc, char* argv[])
      HIPCHECK(hipMalloc(recvbuff+i, recvsize*sizeof(float)));
      h_recvbuff[i] = (float*) malloc (recvsize *sizeof(float));
      HIPCHECK(hipMemset(recvbuff[i], 0, recvsize*sizeof(float)));
-
+      HIPCHECK(hipStreamSynchronize(NULL));
      HIPCHECK(hipStreamCreate(s+i));
  }

@@ -616,7 +616,7 @@ int main(int argc,char* argv[])
      HIPCHECK(hipStreamSynchronize(stream[i]));
      HIPCHECK(hipMemset(d_profiling_data[i], 0, sizeof(struct profiling_data_t)));
    }
-
+    HIPCHECK(hipStreamSynchronize(NULL));
    auto start = std::chrono::high_resolution_clock::now();
    for (int j = 0; j < iters; j ++) {
      for (int i = 0; i < nGpu; i ++) {