AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

SWDEV-259566 - Adding support to retrieve handle for offsetted pointer.

Sfoglia il codice sorgente 
Change-Id: I5da5ab6a24c6df915950637feb486b8c288e60e0
This commit is contained in:
kjayapra-amd
2020-11-10 14:57:02 -05:00
committato da Karthik Jayaprakash
parent 80255c5b26
commit b4492a3de3
4 ha cambiato i file con 65 aggiunte e 42 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
rocclr/hip_internal.hpp
+2 -1
Vedi File
@@ -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;
rocclr/hip_memory.cpp
+5 -3
Vedi File
@@ -2034,7 +2034,7 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* dev_ptr) {
device = hip::getCurrentDevice()->devices()[0];
ihandle = reinterpret_cast<ihipIpcMemHandle_t *>(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);
}
tests/src/ipc/hipSimpleIpc.cpp
+40 -29
Vedi File
@@ -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<void**>(&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<void**>(&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: "<<idx<<std::endl;
}
}
//Close All Mem Handle
HIPCHECK(hipIpcCloseMemHandle(ipc_out_dptr));
HIPCHECK(hipFree(ipc_dptr));
HIPCHECK_RETURN_ONFAIL(hipIpcCloseMemHandle(ipc_out_dptr));
HIPCHECK_RETURN_ONFAIL(hipFree(ipc_dptr));
delete[] ipc_hptr;
delete[] ipc_out_hptr;
return hipSuccess;
}
void multi_process() {
//To create and open IPC handle via multiple process
void positive_cases() {
HIPCHECK(single_process(0));
HIPCHECK(single_process(32));
HIPCHECK(single_process(128));
HIPCHECK(single_process(256));
HIPCHECK(single_process(512));
HIPCHECK(single_process(1023));
HIPCHECK(single_process(47));
HIPCHECK(single_process(191));
HIPCHECK(single_process(1022));
}
void negative_cases() {
HIPCHECK_API(single_process(-1), hipErrorInvalidDevicePointer);
HIPCHECK_API(single_process(1024), hipErrorInvalidDevicePointer);
}
int main() {
single_process();
multi_process();
positive_cases();
negative_cases();
passed();
}
tests/src/test_common.h
+18 -9
Vedi File
@@ -22,6 +22,7 @@ THE SOFTWARE.
* should be added into CPP section
*
*/
#pragma once
#ifdef __cplusplus
#include <iostream>
@@ -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 <tchar.h>
#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<typename T> // 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<typename T>
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<height; i++ ){
for(int j=0; j<width; j++ ){
int offset = i*width + j;
if( input[offset] != output[offset] ){
std::cerr << '[' << i << ',' << j << ',' << "]:" << input[offset] << "----" << output[offset]<<" ";
std::cerr << '[' << i << ',' << j << ',' << "]:" << input[offset]
<< "----" << output[offset]<<" ";
failed("mistmatch at:%d %d",i,j);
}
}
@@ -294,13 +303,13 @@ void initArraysForHost(T** A_h, T** B_h, T** C_h, size_t N, bool usePinnedHost =
if (usePinnedHost) {
if (A_h) {
HIPCHECK(hipHostMalloc((void**)A_h, Nbytes));
HIPCHECK(hipHostMalloc(reinterpret_cast<void**>(A_h), Nbytes));
}
if (B_h) {
HIPCHECK(hipHostMalloc((void**)B_h, Nbytes));
HIPCHECK(hipHostMalloc(reinterpret_cast<void**>(B_h), Nbytes));
}
if (C_h) {
HIPCHECK(hipHostMalloc((void**)C_h, Nbytes));
HIPCHECK(hipHostMalloc(reinterpret_cast<void**>(C_h), Nbytes));
}
} else {
if (A_h) {