From ea701777d8fa8716b77547c150508d79574adb73 Mon Sep 17 00:00:00 2001
From: Vladislav Sytchenko <vladislav.sytchenko@amd.com>
Date: Thu, 19 Mar 2020 16:05:06 -0400
Subject: [PATCH] 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
---
 vdi/hip_texture.cpp | 64 ++++++++++++++++++++++++++++++---------------
 1 file changed, 43 insertions(+), 21 deletions(-)

diff --git a/vdi/hip_texture.cpp b/vdi/hip_texture.cpp
index 2d4085d928..b2e58a7f69 100644
--- a/vdi/hip_texture.cpp
+++ b/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));