diff --git a/rocclr/hip_internal.hpp b/rocclr/hip_internal.hpp index 6ea672666b..248b127d4b 100755 --- a/rocclr/hip_internal.hpp +++ b/rocclr/hip_internal.hpp @@ -48,11 +48,12 @@ /*! IHIP IPC MEMORY Structure */ #define IHIP_IPC_MEM_HANDLE_SIZE 32 -#define IHIP_IPC_MEM_RESERVED_SIZE LP64_SWITCH(28,24) +#define IHIP_IPC_MEM_RESERVED_SIZE LP64_SWITCH(24,16) typedef struct ihipIpcMemHandle_st { char ipc_handle[IHIP_IPC_MEM_HANDLE_SIZE]; ///< ipc memory handle on ROCr size_t psize; + size_t poffset; char reserved[IHIP_IPC_MEM_RESERVED_SIZE]; } ihipIpcMemHandle_t; diff --git a/rocclr/hip_memory.cpp b/rocclr/hip_memory.cpp index dbbbbb0ad0..190d7c6d99 100755 --- a/rocclr/hip_memory.cpp +++ b/rocclr/hip_memory.cpp @@ -2034,7 +2034,7 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* dev_ptr) { device = hip::getCurrentDevice()->devices()[0]; ihandle = reinterpret_cast(handle); - if(!device->IpcCreate(dev_ptr, &(ihandle->psize), &(ihandle->ipc_handle))) { + if(!device->IpcCreate(dev_ptr, &(ihandle->psize), &(ihandle->ipc_handle), &(ihandle->poffset))) { LogPrintfError("IPC memory creation failed for memory: 0x%x", dev_ptr); HIP_RETURN(hipErrorInvalidDevicePointer); } @@ -2061,8 +2061,10 @@ hipError_t hipIpcOpenMemHandle(void** dev_ptr, hipIpcMemHandle_t handle, unsigne HIP_RETURN(hipErrorInvalidValue); } - if(!device->IpcAttach(&(ihandle->ipc_handle), ihandle->psize, flags, dev_ptr)) { - LogPrintfError("cannot attach ipc_handle: with ipc_size: %u flags: %u", ihandle->psize, flags); + if(!device->IpcAttach(&(ihandle->ipc_handle), ihandle->psize, + ihandle->poffset, flags, dev_ptr)) { + LogPrintfError("Cannot attach ipc_handle: with ipc_size: %u" + "ipc_offset: %u flags: %u", ihandle->psize, flags); HIP_RETURN(hipErrorInvalidDevicePointer); } diff --git a/tests/src/ipc/hipSimpleIpc.cpp b/tests/src/ipc/hipSimpleIpc.cpp index 7a4f2c0140..c9c3540dfe 100755 --- a/tests/src/ipc/hipSimpleIpc.cpp +++ b/tests/src/ipc/hipSimpleIpc.cpp @@ -25,67 +25,78 @@ THE SOFTWARE. #include "test_common.h" -#define N 1024 -#define OFFSET 128 +#define NUM 1024 -void single_process() { +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_handle; hipIpcMemHandle_t ipc_offset_handle; - HIPCHECK(hipMalloc((void**)&ipc_dptr, N * sizeof(int))); + HIPCHECK_RETURN_ONFAIL(hipMalloc(reinterpret_cast(&ipc_dptr), NUM * sizeof(int))); - // Negative, Make sure we return error when an offset of original ptr is passed - ipc_offset_dptr = ipc_dptr + (OFFSET * sizeof(int)); - // HIP API return value differs from CUDA's return type - assert(hipErrorInvalidDevicePointer == hipIpcGetMemHandle(&ipc_offset_handle, ipc_offset_dptr)); - - // Get handle for the device_ptr - 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_RETURN_ONFAIL(hipIpcGetMemHandle(&ipc_offset_handle, ipc_offset_dptr)); // Set Values @ Host Ptr - ipc_hptr = new int[N]; - for (size_t idx = 0; idx < N; ++idx) { + ipc_hptr = new int[NUM]; + for (size_t idx = 0; idx < NUM; ++idx) { ipc_hptr[idx] = idx; } // Copy values to Device ptr - HIPCHECK(hipMemset(ipc_dptr, 0x00, (N * sizeof(int)))); - HIPCHECK(hipMemcpy(ipc_dptr, ipc_hptr, (N * sizeof(int)), hipMemcpyHostToDevice)); + 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[N]; - memset(ipc_out_hptr, 0x00, (N * sizeof(int))); - HIPCHECK(hipIpcOpenMemHandle((void**)&ipc_out_dptr, ipc_handle, 0)); + 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(hipMemcpy(ipc_out_hptr, ipc_out_dptr, (N * sizeof(int)), hipMemcpyDeviceToHost)); - for (size_t idx = 0; idx < N; ++idx) { - if(ipc_out_hptr[idx] != idx) { + HIPCHECK_RETURN_ONFAIL(hipMemcpy(ipc_out_hptr, ipc_out_dptr, (NUM * sizeof(int)), hipMemcpyDeviceToHost)); + for (size_t idx = offset; idx < NUM; ++idx) { + if (ipc_out_hptr[idx-offset] != ipc_dptr[idx]) { std::cout<<"Failing @ idx: "< @@ -116,6 +117,14 @@ inline int hip_skip_retcode() { } \ } +#define HIPCHECK_RETURN_ONFAIL(func) \ + do { \ + hipError_t herror = (func); \ + if (herror != hipSuccess) { \ + return herror; \ + } \ + } while (0); + #ifdef _WIN64 #include #define aligned_alloc(x,y) _aligned_malloc(y,x) @@ -185,14 +194,14 @@ int parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg); unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N); template // pointer type -void checkArray(T hData, T hOutputData, size_t width, size_t height,size_t depth) -{ +void checkArray(T hData, T hOutputData, size_t width, size_t height,size_t depth) { for (int i = 0; i < depth; i++) { for (int j = 0; j < height; j++) { for (int k = 0; k < width; k++) { int offset = i*width*height + j*width + k; if (hData[offset] != hOutputData[offset]) { - std::cerr << '[' << i << ',' << j << ',' << k << "]:" << hData[offset] << "----" << hOutputData[offset]<<" "; + std::cerr << '[' << i << ',' << j << ',' << k << "]:" << hData[offset] << "----" + << hOutputData[offset]<<" "; failed("mistmatch at:%d %d %d",i,j,k); } } @@ -201,13 +210,13 @@ void checkArray(T hData, T hOutputData, size_t width, size_t height,size_t depth } template -void checkArray(T input, T output, size_t height, size_t width) -{ +void checkArray(T input, T output, size_t height, size_t width) { for(int i=0; i(A_h), Nbytes)); } if (B_h) { - HIPCHECK(hipHostMalloc((void**)B_h, Nbytes)); + HIPCHECK(hipHostMalloc(reinterpret_cast(B_h), Nbytes)); } if (C_h) { - HIPCHECK(hipHostMalloc((void**)C_h, Nbytes)); + HIPCHECK(hipHostMalloc(reinterpret_cast(C_h), Nbytes)); } } else { if (A_h) {