Support creating textures from user ptr with offset

Change-Id: Ied0af9a842fc6ef067e068a9d89dde63ff5dfce1

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);
 

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)
 {

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));
+                            buffer);
-    // TODO take care of non-zero offset.
+    buffer->release();
-    assert(offset == 0);
     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));
+                            buffer);
-    // TODO take care of non-zero offset.
+    buffer->release();
-    assert(offset == 0);
     if (image == nullptr) {
       return hipErrorInvalidValue;
     }