diff --git a/catch/multiproc/hipIpcMemAccessTest.cc b/catch/multiproc/hipIpcMemAccessTest.cc index 4e143111c4..183da51c24 100644 --- a/catch/multiproc/hipIpcMemAccessTest.cc +++ b/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(&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(&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(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(A_h, C_h, N); - } HIP_CHECK(hipIpcCloseMemHandle(reinterpret_cast(B_d))); HIP_CHECK(hipFree(C_d)); + } } if ((sem_post(sem_ob2)) == -1) { shrd_mem->IfTestPassed = false; diff --git a/catch/unit/memory/hipMemcpy3DAsync.cc b/catch/unit/memory/hipMemcpy3DAsync.cc index 6ebc480844..5918a287c9 100644 --- a/catch/unit/memory/hipMemcpy3DAsync.cc +++ b/catch/unit/memory/hipMemcpy3DAsync.cc @@ -481,58 +481,64 @@ template void Memcpy3DAsync::D2D_SameDeviceMem_StreamDiffDevice() { HIP_CHECK(hipSetDevice(0)); // Allocating the Memory - AllocateMemory(); - HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipStreamCreate(&stream)); - memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); - SetDefaultData(); + int peerAccess = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 0, 1)); + if (peerAccess) { + AllocateMemory(); + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipStreamCreate(&stream)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); - // Host to Device - myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); - myparms.dstArray = arr; + // Host to Device + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); + 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); - HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); - // Array to Array - memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); - SetDefaultData(); - myparms.srcArray = arr; - myparms.dstArray = arr1; + // Array to Array + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + myparms.srcArray = arr; + 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); - HIP_CHECK(hipStreamSynchronize(stream)); - T *hOutputData = reinterpret_cast(malloc(size)); - memset(hOutputData, 0, size); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + T *hOutputData = reinterpret_cast(malloc(size)); + memset(hOutputData, 0, size); - // Device to host - memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); - SetDefaultData(); - myparms.dstPtr = make_hipPitchedPtr(hOutputData, - width * sizeof(T), width, height); - myparms.srcArray = arr1; + // Device to host + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + myparms.dstPtr = make_hipPitchedPtr(hOutputData, + width * sizeof(T), width, height); + 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); - HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); - // Validating the result - HipTest::checkArray(hData, hOutputData, width, height, depth); + // Validating the result + HipTest::checkArray(hData, hOutputData, width, height, depth); - // Deallocating the memory - free(hOutputData); - DeAllocateMemory(); + // Deallocating the memory + free(hOutputData); + DeAllocateMemory(); + } else { + SUCCEED("Skipped the test as there is no peer access"); + } } /* @@ -630,7 +636,7 @@ void Memcpy3DAsync::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;