diff --git a/projects/clr/hipamd/tests/src/ipc/hipMultiProcIpcMem.cpp b/projects/clr/hipamd/tests/src/ipc/hipMultiProcIpcMem.cpp index 15693bfbd6..9301697aa0 100755 --- a/projects/clr/hipamd/tests/src/ipc/hipMultiProcIpcMem.cpp +++ b/projects/clr/hipamd/tests/src/ipc/hipMultiProcIpcMem.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia + * BUILD: %t %s ../test_common.cpp * TEST: %t * HIT_END */ @@ -37,6 +37,7 @@ void multi_process(int num_process, bool debug_process) { int* ipc_hptr = nullptr; int* ipc_out_dptr = nullptr; int* ipc_out_hptr = nullptr; + int* ipc_offset_dptr = nullptr; MultiProcess* mProcess = new MultiProcess(num_process); mProcess->CreateShmem(); @@ -48,7 +49,10 @@ void multi_process(int num_process, bool debug_process) { memset(&ipc_handle, 0x00, sizeof(hipIpcMemHandle_t)); HIPCHECK(hipMalloc((void**)&ipc_dptr, NUM_ELEMS * sizeof(int))); - HIPCHECK(hipIpcGetMemHandle(&ipc_handle, ipc_dptr)); + // Add offset to the dev_ptr + ipc_offset_dptr = ipc_dptr + OFFSET; + // Get handle for the offsetted device_ptr + HIPCHECK(hipIpcGetMemHandle(&ipc_handle, ipc_offset_dptr)); ipc_hptr = new int[NUM_ELEMS]; for (size_t idx = 0; idx < NUM_ELEMS; ++idx) { @@ -68,7 +72,8 @@ void multi_process(int num_process, bool debug_process) { hipIpcMemHandle_t ipc_handle; mProcess->ReadHandleFromShmem(ipc_handle); - HIPCHECK(hipIpcOpenMemHandle((void**)&ipc_out_dptr, ipc_handle, 0)); + // Open handle to get dev_ptr + HIPCHECK(hipIpcOpenMemHandle((void**)&ipc_out_dptr, ipc_handle, hipIpcMemLazyEnablePeerAccess)); HIPCHECK(hipMemcpy(ipc_out_hptr, ipc_out_dptr, (NUM_ELEMS * sizeof(int)), hipMemcpyDeviceToHost)); diff --git a/projects/clr/hipamd/tests/src/ipc/hipSimpleIpcMem.cpp b/projects/clr/hipamd/tests/src/ipc/hipSimpleIpcMem.cpp deleted file mode 100755 index e98201daaa..0000000000 --- a/projects/clr/hipamd/tests/src/ipc/hipSimpleIpcMem.cpp +++ /dev/null @@ -1,102 +0,0 @@ -/* -Copyright (c) 2015-2017 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. -*/ - -/* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia - * TEST: %t - * HIT_END - */ - -#include "test_common.h" - -#define NUM 1024 - -hipError_t single_process(int32_t offset) { - int* ipc_dptr = nullptr; - int* ipc_hptr = nullptr; - int* ipc_out_dptr = nullptr; - int* ipc_out_hptr = nullptr; - - int* ipc_offset_dptr = nullptr; - hipIpcMemHandle_t ipc_offset_handle; - - HIPCHECK_RETURN_ONFAIL(hipMalloc(reinterpret_cast(&ipc_dptr), NUM * sizeof(int))); - - // Add offset to the dev_ptr - ipc_offset_dptr = ipc_dptr + offset; - // Get handle for the offsetted device_ptr - HIPCHECK_RETURN_ONFAIL(hipIpcGetMemHandle(&ipc_offset_handle, ipc_offset_dptr)); - - // Set Values @ Host Ptr - ipc_hptr = new int[NUM]; - for (size_t idx = 0; idx < NUM; ++idx) { - ipc_hptr[idx] = idx; - } - - // Copy values to Device ptr - HIPCHECK_RETURN_ONFAIL(hipMemset(ipc_dptr, 0x00, (NUM * sizeof(int)))); - HIPCHECK_RETURN_ONFAIL(hipMemcpy(ipc_dptr, ipc_hptr, (NUM * sizeof(int)), hipMemcpyHostToDevice)); - - // Open handle to get dev_ptr - ipc_out_hptr = new int[NUM]; - memset(ipc_out_hptr, 0x00, (NUM * sizeof(int))); - HIPCHECK_RETURN_ONFAIL(hipIpcOpenMemHandle(reinterpret_cast(&ipc_out_dptr), - ipc_offset_handle, 0)); - - // Copy Values from Device to Host and Check for correctness - HIPCHECK_RETURN_ONFAIL(hipMemcpy(ipc_out_hptr, ipc_out_dptr, ((NUM-offset) * sizeof(int)), hipMemcpyDeviceToHost)); - for (size_t idx = offset; idx < NUM; ++idx) { - if (ipc_out_hptr[idx-offset] != ipc_dptr[idx]) { - std::cout<<"Failing @ idx: "<