Add 'projects/hip-tests/' from commit '5ce508401e1934053b127de5bf756dcd56a326a2'

git-subtree-dir: projects/hip-tests
git-subtree-mainline: cc7a96c46f
git-subtree-split: 5ce508401e

projects/hip-tests/catch/unit/memory/hipMemcpyWithStreamMultiThread.cc

@@ -0,0 +1,642 @@
+/*
+Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/*
+ * Different test for checking functionality of
+ * hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
+ * hipMemcpyKind kind, hipStream_t stream);
+ */
+
+
+#include <hip_test_common.hh>
+#include <hip_test_kernels.hh>
+#include <vector>
+
+#define LEN 64
+#define SIZE LEN << 2
+#define THREADS 2
+#define MAX_THREADS 16
+
+static constexpr auto N{1024};
+static constexpr auto Nbytes{N * sizeof(int)};
+static const auto MaxGPUDevices{256};
+static constexpr unsigned blocksPerCU{6};  // to hide latency
+static constexpr unsigned threadsPerBlock{256};
+
+enum class ops {
+  TestwithOnestream,
+  TestwithTwoStream,
+  TestOnMultiGPUwithOneStream,
+  TestkindDtoH,
+  TestkindDtoD,
+  TestkindHtoH,
+  TestkindDefault,
+  TestkindDefaultForDtoD,
+  TestDtoDonSameDevice,
+  END_OF_LIST
+};
+
+namespace MemcpyStream {
+void setNumBlocks(int blocksPerCU, int threadsPerBlock, size_t N, unsigned& blocks) {
+  int device;
+  HIP_CHECK_THREAD(hipGetDevice(&device));
+  hipDeviceProp_t props;
+  HIP_CHECK_THREAD(hipGetDeviceProperties(&props, device));
+
+  blocks = props.multiProcessorCount * blocksPerCU;
+  if (blocks * threadsPerBlock > N) {
+    blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
+  }
+}
+}  // namespace MemcpyStream
+
+
+class HipMemcpyWithStreamMultiThreadtests {
+ public:
+  // Test hipMemcpyWithStream with one streams and launch kernel in
+  // that stream, verify the data.
+  void TestwithOnestream(bool& val_res);
+  // Test hipMemcpyWithStream with two streams and launch kernels in
+  // two streams, verify the data.
+  void TestwithTwoStream(bool& val_res);
+  // Test hipMemcpyWithStream with one stream for each gpu and launch
+  // kernels in each, verify the data
+  void TestOnMultiGPUwithOneStream(bool& val_res);
+  // Test hipMemcpyWithStream to copy data from
+  // device to host (hipMemcpyDeviceToHost).
+  void TestkindDtoH(bool& val_res);
+  // Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on MultiGPU.
+  void TestkindDtoD(bool& val_res);
+  // Test hipMemcpyWithStream with hipMemcpyHostToHost.
+  void TestkindHtoH(bool& val_res);
+  // Test hipMemcpyWithStream with hipMemcpyDefault.
+  void TestkindDefault(bool& val_res);
+  // Test hipMemcpyWithStream with hipMemcpyDefault for
+  // device to device transfer case.
+  void TestkindDefaultForDtoD(bool& val_res);
+  // Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on same device.
+  void TestDtoDonSameDevice(bool& val_res);
+  // Allocate Memory
+  void AllocateMemory(int** A_d, int** B_d, int** C_d, int** A_h, int** B_h, int** C_h);
+  // DeAllocate Memory
+  void DeAllocateMemory(int* A_d, int* B_d, int* C_d, int* A_h, int* B_h, int* C_h);
+  // Validate Result
+  bool ValidateResult(int* A_h, int* B_h, int* C_h);
+};
+
+void HipMemcpyWithStreamMultiThreadtests::AllocateMemory(int** A_d, int** B_d, int** C_d, int** A_h,
+                                                         int** B_h, int** C_h) {
+  HIP_CHECK_THREAD(hipMalloc(A_d, Nbytes));
+  HIP_CHECK_THREAD(hipMalloc(B_d, Nbytes));
+  HIP_CHECK_THREAD(hipMalloc(C_d, Nbytes));
+  *A_h = reinterpret_cast<int*>(malloc(Nbytes));
+  *B_h = reinterpret_cast<int*>(malloc(Nbytes));
+  *C_h = reinterpret_cast<int*>(malloc(Nbytes));
+
+  for (size_t i = 0; i < N; i++) {
+    if (*A_h) (*A_h)[i] = 3;
+    if (*B_h) (*B_h)[i] = 4;
+    if (*C_h) (*C_h)[i] = 5;
+  }
+}
+
+void HipMemcpyWithStreamMultiThreadtests::DeAllocateMemory(int* A_d, int* B_d, int* C_d, int* A_h,
+                                                           int* B_h, int* C_h) {
+  HIP_CHECK_THREAD(hipFree(A_d));
+  HIP_CHECK_THREAD(hipFree(B_d));
+  HIP_CHECK_THREAD(hipFree(C_d));
+  free(A_h);
+  free(B_h);
+  free(C_h);
+}
+
+
+bool HipMemcpyWithStreamMultiThreadtests::ValidateResult(int* A_h, int* B_h, int* C_h) {
+  bool TestPassed = true;
+  for (size_t i = 0; i < N; i++) {
+    if ((A_h[i] + B_h[i]) != C_h[i]) {
+      TestPassed = false;
+      break;
+    }
+  }
+  return TestPassed;
+}
+
+
+void HipMemcpyWithStreamMultiThreadtests::TestwithOnestream(bool& val_res) {
+  int *A_d, *B_d, *C_d;
+  int *A_h, *B_h, *C_h;
+  size_t Nbytes{N * sizeof(int)};
+  AllocateMemory(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h);
+  unsigned blocks = 0;
+  MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
+  hipStream_t stream;
+  HIP_CHECK_THREAD(hipStreamCreate(&stream));
+
+
+  HIP_CHECK_THREAD(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
+  HIP_CHECK_THREAD(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream));
+  hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
+                     static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
+  HIP_CHECK_THREAD(hipGetLastError());
+  HIP_CHECK_THREAD(hipStreamSynchronize(stream));
+  HIP_CHECK_THREAD(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+  val_res = ValidateResult(A_h, B_h, C_h);
+  DeAllocateMemory(A_d, B_d, C_d, A_h, B_h, C_h);
+  HIP_CHECK_THREAD(hipStreamDestroy(stream));
+}
+
+void HipMemcpyWithStreamMultiThreadtests::TestwithTwoStream(bool& val_res) {
+  size_t Nbytes = N * sizeof(int);
+  const int NoofStreams = 2;
+  int *A_d[NoofStreams], *B_d[NoofStreams], *C_d[NoofStreams];
+  int *A_h[NoofStreams], *B_h[NoofStreams], *C_h[NoofStreams];
+
+  unsigned blocks = 0;
+  MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
+
+  for (int i = 0; i < NoofStreams; ++i) {
+    AllocateMemory(&A_d[i], &B_d[i], &C_d[i], &A_h[i], &B_h[i], &C_h[i]);
+  }
+
+  hipStream_t stream[NoofStreams];
+  for (int i = 0; i < NoofStreams; ++i) {
+    HIP_CHECK_THREAD(hipStreamCreate(&stream[i]));
+  }
+
+  for (int i = 0; i < NoofStreams; ++i) {
+    HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
+    HIP_CHECK_THREAD(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
+  }
+
+  for (int i = 0; i < NoofStreams; ++i) {
+    hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
+                       static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
+    HIP_CHECK_THREAD(hipGetLastError());
+  }
+
+  for (int i = 0; i < NoofStreams; ++i) {
+    HIP_CHECK_THREAD(hipStreamSynchronize(stream[i]));
+    HIP_CHECK_THREAD(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
+    val_res = ValidateResult(A_h[i], B_h[i], C_h[i]);
+  }
+
+  for (int i = 0; i < NoofStreams; ++i) {
+    DeAllocateMemory(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i]);
+    HIP_CHECK_THREAD(hipStreamDestroy(stream[i]));
+  }
+}
+
+void HipMemcpyWithStreamMultiThreadtests::TestDtoDonSameDevice(bool& val_res) {
+  size_t Nbytes = N * sizeof(int);
+  const int NoofStreams = 2;
+  int *A_d[NoofStreams], *B_d[NoofStreams], *C_d[NoofStreams];
+  int *A_h[NoofStreams], *B_h[NoofStreams], *C_h[NoofStreams];
+
+  unsigned blocks = 0;
+  MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
+
+  AllocateMemory(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0]);
+
+
+  hipStream_t stream[NoofStreams];
+  for (int i = 0; i < NoofStreams; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(0));
+    HIP_CHECK_THREAD(hipStreamCreate(&stream[i]));
+  }
+
+  HIP_CHECK_THREAD(hipSetDevice(0));
+  HIP_CHECK_THREAD(hipMalloc(&A_d[1], Nbytes));
+  HIP_CHECK_THREAD(hipMalloc(&B_d[1], Nbytes));
+  HIP_CHECK_THREAD(hipMalloc(&C_d[1], Nbytes));
+  C_h[1] = reinterpret_cast<int*>(malloc(Nbytes));
+  HIPASSERT(C_h[1] != NULL);
+
+  HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
+  HIP_CHECK_THREAD(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
+
+  HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[1], A_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[1]));
+  HIP_CHECK_THREAD(hipMemcpyWithStream(B_d[1], B_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[1]));
+
+
+  for (int i = 0; i < NoofStreams; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(0));
+    hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
+                       static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
+    HIP_CHECK_THREAD(hipGetLastError());
+  }
+
+  for (int i = 0; i < NoofStreams; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(0));
+    HIP_CHECK_THREAD(hipStreamSynchronize(stream[i]));
+    HIP_CHECK_THREAD(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
+    val_res = ValidateResult(A_h[0], B_h[0], C_h[i]);
+  }
+
+
+  DeAllocateMemory(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0]);
+
+  if (A_d[1]) {
+    HIP_CHECK_THREAD(hipFree(A_d[1]));
+  }
+  if (B_d[1]) {
+    HIP_CHECK_THREAD(hipFree(B_d[1]));
+  }
+  if (C_d[1]) {
+    HIP_CHECK_THREAD(hipFree(C_d[1]));
+  }
+  if (C_h[1]) {
+    free(C_h[1]);
+  }
+
+
+  for (int i = 0; i < NoofStreams; ++i) {
+    HIP_CHECK_THREAD(hipStreamDestroy(stream[i]));
+  }
+}
+
+void HipMemcpyWithStreamMultiThreadtests::TestOnMultiGPUwithOneStream(bool& val_res) {
+  size_t Nbytes = N * sizeof(int);
+  int numDevices = 0;
+
+  unsigned blocks = 0;
+  MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
+  HIP_CHECK_THREAD(hipGetDeviceCount(&numDevices));
+  // If you have single GPU machine the return
+  if (numDevices <= 1) {
+    return;
+  }
+  int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices];
+  int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices];
+
+  hipStream_t stream[MaxGPUDevices];
+  for (int i = 0; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i));
+    HIP_CHECK_THREAD(hipStreamCreate(&stream[i]));
+  }
+
+  for (int i = 0; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i));
+    AllocateMemory(&A_d[i], &B_d[i], &C_d[i], &A_h[i], &B_h[i], &C_h[i]);
+  }
+
+
+  for (int i = 0; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i));
+    HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
+    HIP_CHECK_THREAD(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
+  }
+
+  for (int i = 0; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i));
+    hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
+                       static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
+    HIP_CHECK_THREAD(hipGetLastError());
+  }
+
+  for (int i = 0; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i));
+    HIP_CHECK_THREAD(hipStreamSynchronize(stream[i]));
+    HIP_CHECK_THREAD(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
+    val_res = ValidateResult(A_h[i], B_h[i], C_h[i]);
+  }
+
+  for (int i = 0; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i));
+    DeAllocateMemory(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i]);
+    HIP_CHECK_THREAD(hipStreamDestroy(stream[i]));
+  }
+}
+
+void HipMemcpyWithStreamMultiThreadtests::TestkindDtoH(bool& val_res) {
+  size_t Nbytes = N * sizeof(int);
+  int *A_d, *B_d, *C_d;
+  int *A_h, *B_h, *C_h;
+
+  unsigned blocks = 0;
+  MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
+  AllocateMemory(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h);
+
+  hipStream_t stream;
+  HIP_CHECK_THREAD(hipStreamCreate(&stream));
+
+  HIP_CHECK_THREAD(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
+  HIP_CHECK_THREAD(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream));
+  hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
+                     static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
+  HIP_CHECK_THREAD(hipGetLastError());
+  HIP_CHECK_THREAD(hipStreamSynchronize(stream));
+  HIP_CHECK_THREAD(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
+  val_res = ValidateResult(A_h, B_h, C_h);
+
+  DeAllocateMemory(A_d, B_d, C_d, A_h, B_h, C_h);
+  HIP_CHECK_THREAD(hipStreamDestroy(stream));
+}
+
+
+void HipMemcpyWithStreamMultiThreadtests::TestkindDtoD(bool& val_res) {
+  size_t Nbytes = N * sizeof(int);
+  int numDevices = 0;
+
+
+  unsigned blocks = 0;
+  MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
+  HIP_CHECK_THREAD(hipGetDeviceCount(&numDevices));
+  // If you have single GPU machine the return
+  if (numDevices <= 1) {
+    return;
+  }
+
+  int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices];
+  int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices];
+
+  hipStream_t stream[MaxGPUDevices];
+  for (int i = 0; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i));
+    HIP_CHECK_THREAD(hipStreamCreate(&stream[i]));
+  }
+
+  // Initialize and create the host and device elements for first device
+  HIP_CHECK_THREAD(hipSetDevice(0));
+  AllocateMemory(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0]);
+
+  for (int i = 1; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i))
+    HIP_CHECK_THREAD(hipMalloc(&A_d[i], Nbytes));
+    HIP_CHECK_THREAD(hipMalloc(&B_d[i], Nbytes));
+    HIP_CHECK_THREAD(hipMalloc(&C_d[i], Nbytes));
+    C_h[i] = reinterpret_cast<int*>(malloc(Nbytes));
+    HIPASSERT(C_h[i] != NULL);
+  }
+
+
+  HIP_CHECK_THREAD(hipSetDevice(0));
+  HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
+  HIP_CHECK_THREAD(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
+
+  // Copying device data from 1st GPU to the rest of the the GPUs that is
+  // numDevices in the setup. 1st GPU start numbering from 0,1,2..n etc.
+  for (int i = 1; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i));
+    HIP_CHECK_THREAD(
+        hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[i]));
+    HIP_CHECK_THREAD(
+        hipMemcpyWithStream(B_d[i], B_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[i]));
+  }
+
+
+  // Launching the kernel including the 1st GPU to the no of GPUs present
+  // in the setup. 1st GPU start numbering from 0,1,2..n etc.
+  for (int i = 0; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i));
+    hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
+                       static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
+    HIP_CHECK_THREAD(hipGetLastError());
+  }
+
+  for (int i = 0; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i));
+    HIP_CHECK_THREAD(hipStreamSynchronize(stream[i]));
+    HIP_CHECK_THREAD(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
+    val_res = ValidateResult(A_h[0], B_h[0], C_h[i]);
+  }
+
+  DeAllocateMemory(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0]);
+  HIP_CHECK_THREAD(hipStreamDestroy(stream[0]));
+
+  for (int i = 1; i < numDevices; ++i) {
+    if (A_d[i]) {
+      HIP_CHECK_THREAD(hipFree(A_d[i]));
+    }
+    if (B_d[i]) {
+      HIP_CHECK_THREAD(hipFree(B_d[i]));
+    }
+    if (C_d[i]) {
+      HIP_CHECK_THREAD(hipFree(C_d[i]));
+    }
+    if (C_h[i]) {
+      free(C_h[i]);
+    }
+    HIP_CHECK_THREAD(hipStreamDestroy(stream[i]));
+  }
+}
+
+void HipMemcpyWithStreamMultiThreadtests::TestkindDefault(bool& val_res) {
+  size_t Nbytes = N * sizeof(int);
+  int *A_d, *B_d, *C_d;
+  int *A_h, *B_h, *C_h;
+
+
+  unsigned blocks = 0;
+  MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
+  AllocateMemory(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h);
+
+  hipStream_t stream;
+  HIP_CHECK_THREAD(hipStreamCreate(&stream));
+
+  HIP_CHECK_THREAD(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyDefault, stream));
+  HIP_CHECK_THREAD(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyDefault, stream));
+  hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
+                     static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
+  HIP_CHECK_THREAD(hipGetLastError());
+  HIP_CHECK_THREAD(hipStreamSynchronize(stream));
+  HIP_CHECK_THREAD(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDefault, stream));
+  val_res = ValidateResult(A_h, B_h, C_h);
+
+  DeAllocateMemory(A_d, B_d, C_d, A_h, B_h, C_h);
+  HIP_CHECK_THREAD(hipStreamDestroy(stream));
+}
+
+void HipMemcpyWithStreamMultiThreadtests::TestkindDefaultForDtoD(bool& val_res) {
+  size_t Nbytes = N * sizeof(int);
+  int numDevices = 0;
+
+
+  unsigned blocks = 0;
+  MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
+  HIP_CHECK_THREAD(hipGetDeviceCount(&numDevices));
+  // Test case will not run on single GPU setup.
+  if (numDevices <= 1) {
+    return;
+  }
+
+  int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices];
+  int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices];
+
+  // Initialize and create the host and device elements for first device
+  HIP_CHECK_THREAD(hipSetDevice(0));
+  AllocateMemory(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0]);
+
+  for (int i = 1; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i));
+    HIP_CHECK_THREAD(hipMalloc(&A_d[i], Nbytes));
+    HIP_CHECK_THREAD(hipMalloc(&B_d[i], Nbytes));
+    HIP_CHECK_THREAD(hipMalloc(&C_d[i], Nbytes));
+    C_h[i] = reinterpret_cast<int*>(malloc(Nbytes));
+    HIPASSERT(C_h[i] != NULL);
+  }
+
+  hipStream_t stream[MaxGPUDevices];
+  for (int i = 0; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i));
+    HIP_CHECK_THREAD(hipStreamCreate(&stream[i]));
+  }
+
+  HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
+  HIP_CHECK_THREAD(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
+
+  // Copying device data from 1st GPU to the rest of the the GPUs
+  // using hipMemcpyDefault kind  that is numDevices in the setup.
+  // 1st GPU start numbering from 0,1,2..n etc.
+  for (int i = 1; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, hipMemcpyDefault, stream[i]));
+    HIP_CHECK_THREAD(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes, hipMemcpyDefault, stream[i]));
+  }
+
+  for (int i = 0; i < numDevices; ++i) {
+    hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
+                       static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
+    HIP_CHECK_THREAD(hipGetLastError());
+  }
+
+  for (int i = 0; i < numDevices; ++i) {
+    HIP_CHECK_THREAD(hipSetDevice(i));  // hipMemcpy will be on this device
+    HIP_CHECK_THREAD(hipStreamSynchronize(stream[i]));
+    HIP_CHECK_THREAD(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
+    // Output of each GPU is getting validated with input of 1st GPU.
+    val_res = ValidateResult(A_h[0], B_h[0], C_h[i]);
+  }
+
+  DeAllocateMemory(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0]);
+  HIP_CHECK_THREAD(hipStreamDestroy(stream[0]));
+
+  for (int i = 1; i < numDevices; ++i) {
+    if (A_d[i]) {
+      HIP_CHECK_THREAD(hipFree(A_d[i]));
+    }
+    if (B_d[i]) {
+      HIP_CHECK_THREAD(hipFree(B_d[i]));
+    }
+    if (C_d[i]) {
+      HIP_CHECK_THREAD(hipFree(C_d[i]));
+    }
+    if (C_h[i]) {
+      free(C_h[i]);
+    }
+    HIP_CHECK_THREAD(hipStreamDestroy(stream[i]));
+  }
+}
+
+void HipMemcpyWithStreamMultiThreadtests::TestkindHtoH(bool& val_res) {
+  size_t Nbytes = N * sizeof(int);
+  int *A_h, *B_h;
+
+  // Allocate memory to A_h and B_h
+  A_h = static_cast<int*>(malloc(Nbytes));
+  B_h = static_cast<int*>(malloc(Nbytes));
+
+  for (size_t i = 0; i < N; ++i) {
+    if (A_h) {
+      (A_h)[i] = 3.146f + i;  // Pi
+    }
+  }
+
+  hipStream_t stream;
+  HIP_CHECK_THREAD(hipStreamCreate(&stream));
+
+  HIP_CHECK_THREAD(hipMemcpyWithStream(B_h, A_h, Nbytes, hipMemcpyHostToHost, stream));
+  HIP_CHECK_THREAD(hipStreamSynchronize(stream));
+
+  for (size_t i = 0; i < N; i++) {
+    if ((A_h[i] != B_h[i])) {
+      val_res = false;
+      break;
+    }
+  }
+
+  if (A_h) {
+    free(A_h);
+  }
+  if (B_h) {
+    free(B_h);
+  }
+  HIP_CHECK_THREAD(hipStreamDestroy(stream));
+}
+
+TEST_CASE("Unit_hipMemcpyWithStream_MultiThread") {
+  const auto Threadcount{10};
+  bool ret_val[Threadcount];
+  std::thread th[Threadcount];
+  for (int op = static_cast<int>(ops::TestwithOnestream); op < static_cast<int>(ops::END_OF_LIST);
+       ++op) {
+    HipMemcpyWithStreamMultiThreadtests tests;
+    for (uint32_t i = 0; i < Threadcount; i++) {
+      switch (static_cast<ops>(op)) {
+        case ops::TestwithOnestream:
+          th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestwithOnestream, &tests,
+                              std::ref(ret_val[i]));
+          break;
+        case ops::TestwithTwoStream:
+          th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestwithTwoStream, &tests,
+                              std::ref(ret_val[i]));
+          break;
+        case ops::TestkindDtoH:
+          th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestkindDtoH, &tests,
+                              std::ref(ret_val[i]));
+          break;
+        case ops::TestkindHtoH:
+          th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestkindHtoH, &tests,
+                              std::ref(ret_val[i]));
+          break;
+        case ops::TestkindDtoD:
+          th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestkindDtoD, &tests,
+                              std::ref(ret_val[i]));
+          break;
+        case ops::TestOnMultiGPUwithOneStream:
+          th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestOnMultiGPUwithOneStream,
+                              &tests, std::ref(ret_val[i]));
+          break;
+        case ops::TestkindDefault:
+          th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestkindDefault, &tests,
+                              std::ref(ret_val[i]));
+          break;
+        case ops::TestkindDefaultForDtoD:
+          th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestkindDefaultForDtoD, &tests,
+                              std::ref(ret_val[i]));
+          break;
+        case ops::TestDtoDonSameDevice:
+          th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestDtoDonSameDevice, &tests,
+                              std::ref(ret_val[i]));
+          break;
+        default: {
+        }
+      }
+    }
+
+    for (uint32_t i = 0; i < Threadcount; i++) {
+      th[i].join();
+    }
+
+    HIP_CHECK_THREAD_FINALIZE();
+
+    for (uint32_t i = 0; i < Threadcount; i++) {
+      REQUIRE(ret_val[i] == true);
+    }
+  }
+}