SWDEV-323848 - Fix for p2p failed tests (#2537)

Added peer check for the p2p related testcases

Change-Id: I9131b284baf33bef18eceaa7ed462d733f91be92

catch/multiproc/hipIpcMemAccessTest.cc

@@ -118,12 +118,12 @@ TEST_CASE("Unit_hipIpcMemAccess_Semaphores") {
       }
       for (int i = 0; i < Num_devices; ++i) {
         HIP_CHECK(hipSetDevice(i));
-        HIP_CHECK(hipMalloc(&C_d, Nbytes));
-        HIP_CHECK(hipIpcOpenMemHandle(reinterpret_cast<void **>(&B_d),
-                                      shrd_mem->memHandle,
-                                      hipIpcMemLazyEnablePeerAccess));
         HIP_CHECK(hipDeviceCanAccessPeer(&CanAccessPeer, i, shrd_mem->device));
         if (CanAccessPeer == 1) {
+          HIP_CHECK(hipMalloc(&C_d, Nbytes));
+          HIP_CHECK(hipIpcOpenMemHandle(reinterpret_cast<void **>(&B_d),
+                                        shrd_mem->memHandle,
+                                        hipIpcMemLazyEnablePeerAccess));
           HIP_CHECK(hipMemcpy(C_d, B_d, Nbytes, hipMemcpyDeviceToDevice));
           HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
           HipTest::checkTest<int>(A_h, C_h, N);
@@ -131,9 +131,9 @@ TEST_CASE("Unit_hipIpcMemAccess_Semaphores") {
           // Checking if the data obtained from Ipc shared memory is consistent
           HIP_CHECK(hipMemcpy(C_h, B_d, Nbytes, hipMemcpyDeviceToHost));
           HipTest::checkTest<int>(A_h, C_h, N);
-        }
         HIP_CHECK(hipIpcCloseMemHandle(reinterpret_cast<void*>(B_d)));
         HIP_CHECK(hipFree(C_d));
+        }
       }
       if ((sem_post(sem_ob2)) == -1) {
         shrd_mem->IfTestPassed = false;

catch/unit/memory/hipMemcpy3DAsync.cc

@@ -481,58 +481,64 @@ template <typename T>
 void Memcpy3DAsync<T>::D2D_SameDeviceMem_StreamDiffDevice() {
   HIP_CHECK(hipSetDevice(0));
   // Allocating the Memory
-  AllocateMemory();
+  int peerAccess = 0;
-  HIP_CHECK(hipSetDevice(1));
+  HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 0, 1));
-  HIP_CHECK(hipStreamCreate(&stream));
+  if (peerAccess) {
-  memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
+    AllocateMemory();
-  SetDefaultData();
+    HIP_CHECK(hipSetDevice(1));
+    HIP_CHECK(hipStreamCreate(&stream));
+    memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
+    SetDefaultData();
 
-  // Host to Device
+    // Host to Device
-  myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
+    myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
-  myparms.dstArray = arr;
+    myparms.dstArray = arr;
 #ifdef __HIP_PLATFORM_NVCC__
-  myparms.kind = cudaMemcpyHostToDevice;
+    myparms.kind = cudaMemcpyHostToDevice;
 #else
-  myparms.kind = hipMemcpyHostToDevice;
+    myparms.kind = hipMemcpyHostToDevice;
 #endif
-  REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
+    REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
-  HIP_CHECK(hipStreamSynchronize(stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
 
-  // Array to Array
+    // Array to Array
-  memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
+    memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
-  SetDefaultData();
+    SetDefaultData();
-  myparms.srcArray = arr;
+    myparms.srcArray = arr;
-  myparms.dstArray = arr1;
+    myparms.dstArray = arr1;
 #ifdef __HIP_PLATFORM_NVCC__
-  myparms.kind = cudaMemcpyDeviceToDevice;
+    myparms.kind = cudaMemcpyDeviceToDevice;
 #else
-  myparms.kind = hipMemcpyDeviceToDevice;
+    myparms.kind = hipMemcpyDeviceToDevice;
 #endif
-  REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
+    REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
-  HIP_CHECK(hipStreamSynchronize(stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
-  T *hOutputData = reinterpret_cast<T*>(malloc(size));
+    T *hOutputData = reinterpret_cast<T*>(malloc(size));
-  memset(hOutputData, 0,  size);
+    memset(hOutputData, 0,  size);
 
-  // Device to host
+    // Device to host
-  memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
+    memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
-  SetDefaultData();
+    SetDefaultData();
-  myparms.dstPtr = make_hipPitchedPtr(hOutputData,
+    myparms.dstPtr = make_hipPitchedPtr(hOutputData,
-                   width * sizeof(T), width, height);
+        width * sizeof(T), width, height);
-  myparms.srcArray = arr1;
+    myparms.srcArray = arr1;
 #ifdef __HIP_PLATFORM_NVCC__
-  myparms.kind = cudaMemcpyDeviceToHost;
+    myparms.kind = cudaMemcpyDeviceToHost;
 #else
-  myparms.kind = hipMemcpyDeviceToHost;
+    myparms.kind = hipMemcpyDeviceToHost;
 #endif
-  REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
+    REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
-  HIP_CHECK(hipStreamSynchronize(stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
 
-  // Validating the result
+    // Validating the result
-  HipTest::checkArray(hData, hOutputData, width, height, depth);
+    HipTest::checkArray(hData, hOutputData, width, height, depth);
 
-  // Deallocating the memory
+    // Deallocating the memory
-  free(hOutputData);
+    free(hOutputData);
-  DeAllocateMemory();
+    DeAllocateMemory();
+  } else {
+    SUCCEED("Skipped the test as there is no peer access");
+  }
 }
 
 /*
@@ -630,7 +636,7 @@ void Memcpy3DAsync<T>::simple_Memcpy3DAsync() {
   memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
   SetDefaultData();
   myparms.dstPtr = make_hipPitchedPtr(hOutputData,
-                   width * sizeof(T), width, height);
+      width * sizeof(T), width, height);
   myparms.srcArray = arr1;
 #ifdef __HIP_PLATFORM_NVCC__
   myparms.kind = cudaMemcpyDeviceToHost;