From f8d11f2f30c94da4fca2b0a0e9b37c1f9ef7be72 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


[ROCm/hip commit: ea701777d8fa8716b77547c150508d79574adb73]
---
 projects/hip/vdi/hip_texture.cpp | 64 +++++++++++++++++++++-----------
 1 file changed, 43 insertions(+), 21 deletions(-)

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