diff --git a/hipamd/rocclr/hip_internal.hpp b/hipamd/rocclr/hip_internal.hpp index a53db4be74..de41120f79 100755 --- a/hipamd/rocclr/hip_internal.hpp +++ b/hipamd/rocclr/hip_internal.hpp @@ -236,6 +236,7 @@ extern hipError_t ihipDeviceGetCount(int* count); extern int ihipGetDevice(); extern hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags); extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset); +extern amd::Memory* getMemoryObjectWithOffset(const void* ptr, const size_t size); extern bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, void** var_addr, size_t* var_size); diff --git a/hipamd/rocclr/hip_memory.cpp b/hipamd/rocclr/hip_memory.cpp index a3771a5e20..a8405e595c 100755 --- a/hipamd/rocclr/hip_memory.cpp +++ b/hipamd/rocclr/hip_memory.cpp @@ -43,6 +43,27 @@ amd::Memory* getMemoryObject(const void* ptr, size_t& offset) { return memObj; } +// ================================================================================================ +amd::Memory* getMemoryObjectWithOffset(const void* ptr, const size_t size) { + size_t offset; + amd::Memory* memObj = getMemoryObject(ptr, offset); + + if (memObj != nullptr) { + assert(size <= (memObj->getSize() - offset)); + memObj = new (memObj->getContext()) amd::Buffer(*memObj, memObj->getMemFlags(), offset, size); + if (memObj == nullptr) {; + return nullptr; + } + + if (!memObj->create(nullptr)) { + memObj->release(); + return nullptr; + } + } + + return memObj; +} + // ================================================================================================ hipError_t ihipFree(void *ptr) { diff --git a/hipamd/rocclr/hip_texture.cpp b/hipamd/rocclr/hip_texture.cpp index d56178e4ee..4980f87eaf 100755 --- a/hipamd/rocclr/hip_texture.cpp +++ b/hipamd/rocclr/hip_texture.cpp @@ -247,20 +247,20 @@ hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject, const cl_channel_type channelType = hip::getCLChannelType(hip::getArrayFormat(pResDesc->res.linear.desc), pTexDesc->readMode); const amd::Image::Format imageFormat({channelOrder, channelType}); const cl_mem_object_type imageType = hip::getCLMemObjectType(pResDesc->resType); - size_t offset = 0; + const size_t imageSizeInBytes = pResDesc->res.linear.sizeInBytes; + amd::Memory* buffer = getMemoryObjectWithOffset(pResDesc->res.linear.devPtr, imageSizeInBytes); image = ihipImageCreate(channelOrder, channelType, imageType, - (pResDesc->res.linear.sizeInBytes / imageFormat.getElementSize()), /* imageWidth */ + imageSizeInBytes / imageFormat.getElementSize(), /* imageWidth */ 0, /* imageHeight */ 0, /* imageDepth */ 0, /* imageArraySize */ 0, /* imageRowPitch */ 0, /* imageSlicePitch */ 0, /* numMipLevels */ - getMemoryObject(pResDesc->res.linear.devPtr, offset)); - // TODO take care of non-zero offset. - assert(offset == 0); + buffer); + buffer->release(); if (image == nullptr) { return hipErrorInvalidValue; } @@ -270,7 +270,8 @@ hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject, const cl_channel_order channelOrder = hip::getCLChannelOrder(hip::getNumChannels(pResDesc->res.pitch2D.desc), pTexDesc->sRGB); const cl_channel_type channelType = hip::getCLChannelType(hip::getArrayFormat(pResDesc->res.pitch2D.desc), pTexDesc->readMode); const cl_mem_object_type imageType = hip::getCLMemObjectType(pResDesc->resType); - size_t offset = 0; + const size_t imageSizeInBytes = pResDesc->res.pitch2D.pitchInBytes * pResDesc->res.pitch2D.height; + amd::Memory* buffer = getMemoryObjectWithOffset(pResDesc->res.pitch2D.devPtr, imageSizeInBytes); image = ihipImageCreate(channelOrder, channelType, imageType, @@ -281,9 +282,8 @@ hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject, pResDesc->res.pitch2D.pitchInBytes, /* imageRowPitch */ 0, /* imageSlicePitch */ 0, /* numMipLevels */ - getMemoryObject(pResDesc->res.pitch2D.devPtr, offset)); - // TODO take care of non-zero offset. - assert(offset == 0); + buffer); + buffer->release(); if (image == nullptr) { return hipErrorInvalidValue; }