Allow creating texture from unaligned user ptr

All we have to do is align the ptr to HW requirments an if it's not zero, then return the offset to the user.

We currently don't have anywhere to store this offset, so hipGetTextureAlignmentOffset() will still always return 0.

Change-Id: If31998127d99a2a3222a026d88249519d6102505


[ROCm/hip commit: ea701777d8]

projects/hip/vdi/hip_texture.cpp

@@ -382,6 +382,28 @@ hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc,
   HIP_RETURN(hipSuccess);
 }
 
+inline bool ihipGetTextureAlignmentOffset(size_t* offset,
+                                          const void* devPtr) {
+  amd::Device* device = hip::getCurrentDevice()->devices()[0];
+  const device::Info& info = device->info();
+
+  const char* alignedDevPtr = amd::alignUp(static_cast<const char*>(devPtr), info.imageBaseAddressAlignment_);
+  const size_t alignedOffset = alignedDevPtr - static_cast<const char*>(devPtr);
+
+  // If the device memory pointer was returned from hipMalloc(),
+  // the offset is guaranteed to be 0 and NULL may be passed as the offset parameter.
+  if ((alignedOffset != 0) &&
+      (offset == nullptr)) {
+    return false;
+  }
+
+  if (offset != nullptr) {
+    *offset = alignedOffset;
+  }
+
+  return true;
+}
+
 hipError_t ihipBindTexture(size_t* offset,
                            const textureReference* texref,
                            const void* devPtr,
@@ -398,19 +420,19 @@ hipError_t ihipBindTexture(size_t* offset,
   // No need to check for errors.
   ihipDestroyTextureObject(texref->textureObject);
 
-  // If the device memory pointer was returned from hipMalloc(),
-  // the offset is guaranteed to be 0 and NULL may be passed as the offset parameter.
-  // TODO enforce alignment on devPtr.
-  if (offset != nullptr) {
-    *offset = 0;
-  }
-
   hipResourceDesc resDesc = {};
   resDesc.resType = hipResourceTypeLinear;
   resDesc.res.linear.devPtr = const_cast<void*>(devPtr);
   resDesc.res.linear.desc = *desc;
   resDesc.res.linear.sizeInBytes = size;
 
+  if (ihipGetTextureAlignmentOffset(offset, devPtr)) {
+    // Align the user ptr to HW requirments.
+    resDesc.res.linear.devPtr = static_cast<char*>(const_cast<void*>(devPtr)) - *offset;
+  } else {
+    return hipErrorInvalidValue;
+  }
+
   hipTextureDesc texDesc = hip::getTextureDesc(texref);
 
   return ihipCreateTextureObject(const_cast<hipTextureObject_t*>(&texref->textureObject), &resDesc, &texDesc, nullptr);
@@ -434,13 +456,6 @@ hipError_t ihipBindTexture2D(size_t* offset,
   // No need to check for errors.
   ihipDestroyTextureObject(texref->textureObject);
 
-  // If the device memory pointer was returned from hipMalloc(),
-  // the offset is guaranteed to be 0 and NULL may be passed as the offset parameter.
-  // TODO enforce alignment on devPtr.
-  if (offset != nullptr) {
-    *offset = 0;
-  }
-
   hipResourceDesc resDesc = {};
   resDesc.resType = hipResourceTypePitch2D;
   resDesc.res.pitch2D.devPtr = const_cast<void*>(devPtr);
@@ -449,6 +464,13 @@ hipError_t ihipBindTexture2D(size_t* offset,
   resDesc.res.pitch2D.height = height;
   resDesc.res.pitch2D.pitchInBytes = pitch;
 
+  if (ihipGetTextureAlignmentOffset(offset, devPtr)) {
+    // Align the user ptr to HW requirments.
+    resDesc.res.pitch2D.devPtr = static_cast<char*>(const_cast<void*>(devPtr)) - *offset;
+  } else {
+    return hipErrorInvalidValue;
+  }
+
   hipTextureDesc texDesc = hip::getTextureDesc(texref);
 
   return ihipCreateTextureObject(const_cast<hipTextureObject_t*>(&texref->textureObject), &resDesc, &texDesc, nullptr);
@@ -798,19 +820,19 @@ hipError_t hipTexRefSetAddress(size_t* ByteOffset,
   // No need to check for errors.
   ihipDestroyTextureObject(texRef->textureObject);
 
-  // If the device memory pointer was returned from hipMemAlloc(),
-  // the offset is guaranteed to be 0 and NULL may be passed as the ByteOffset parameter.
-  // TODO enforce alignment on devPtr.
-  if (ByteOffset != nullptr) {
-    *ByteOffset = 0;
-  }
-
   hipResourceDesc resDesc = {};
   resDesc.resType = hipResourceTypeLinear;
   resDesc.res.linear.devPtr = dptr;
   resDesc.res.linear.desc = hip::getChannelFormatDesc(texRef->numChannels, texRef->format);
   resDesc.res.linear.sizeInBytes = bytes;
 
+  if (ihipGetTextureAlignmentOffset(ByteOffset, dptr)) {
+    // Align the user ptr to HW requirments.
+    resDesc.res.linear.devPtr = static_cast<char*>(dptr) - *ByteOffset;
+  } else {
+    return HIP_RETURN(hipErrorInvalidValue);
+  }
+
   hipTextureDesc texDesc = hip::getTextureDesc(texRef);
 
   HIP_RETURN(ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, nullptr));