From cf49aaa7fcd744ef774dd397398fa76cd0264aa4 Mon Sep 17 00:00:00 2001
From: Vladislav Sytchenko <vladislav.sytchenko@amd.com>
Date: Mon, 23 Mar 2020 15:53:29 -0400
Subject: [PATCH] Replace hip::TextureObject with __hip_texture

This avoids the use of extra casts when obtaining a texture object handle.

Change-Id: I42df22bdad0ab9ac6c33cb8b282dee65fe7cfd6e


[ROCm/clr commit: 6cfbe19160706b83b14e5c246488036eb9f6c742]
---
 projects/clr/hipamd/vdi/hip_texture.cpp | 74 ++++++++++++-------------
 1 file changed, 34 insertions(+), 40 deletions(-)

diff --git a/projects/clr/hipamd/vdi/hip_texture.cpp b/projects/clr/hipamd/vdi/hip_texture.cpp
index 19761d1e78..f838d6cb0c 100644
--- a/projects/clr/hipamd/vdi/hip_texture.cpp
+++ b/projects/clr/hipamd/vdi/hip_texture.cpp
@@ -24,36 +24,34 @@
 #include "hip_conversions.hpp"
 #include "platform/sampler.hpp"
 
-namespace hip {
-  struct TextureObject {
-    uint32_t imageSRD[HIP_IMAGE_OBJECT_SIZE_DWORD];
-    uint32_t samplerSRD[HIP_SAMPLER_OBJECT_SIZE_DWORD];
-    amd::Image* image;
-    amd::Sampler* sampler;
-    hipResourceDesc resDesc;
-    hipTextureDesc texDesc;
-    hipResourceViewDesc resViewDesc;
+struct __hip_texture {
+  uint32_t imageSRD[HIP_IMAGE_OBJECT_SIZE_DWORD];
+  uint32_t samplerSRD[HIP_SAMPLER_OBJECT_SIZE_DWORD];
+  amd::Image* image;
+  amd::Sampler* sampler;
+  hipResourceDesc resDesc;
+  hipTextureDesc texDesc;
+  hipResourceViewDesc resViewDesc;
 
-    TextureObject(amd::Image* image_,
-                  amd::Sampler* sampler_,
-                  const hipResourceDesc& resDesc_,
-                  const hipTextureDesc& texDesc_,
-                  const hipResourceViewDesc& resViewDesc_) :
-      image(image_),
-      sampler(sampler_),
-      resDesc(resDesc_),
-      texDesc(texDesc_),
-      resViewDesc(resViewDesc_) {
-      amd::Context& context = *hip::getCurrentDevice()->asContext();
-      amd::Device& device = *context.devices()[0];
+  __hip_texture(amd::Image* image_,
+                amd::Sampler* sampler_,
+                const hipResourceDesc& resDesc_,
+                const hipTextureDesc& texDesc_,
+                const hipResourceViewDesc& resViewDesc_) :
+    image(image_),
+    sampler(sampler_),
+    resDesc(resDesc_),
+    texDesc(texDesc_),
+    resViewDesc(resViewDesc_) {
+    amd::Context& context = *hip::getCurrentDevice()->asContext();
+    amd::Device& device = *context.devices()[0];
 
-      device::Memory* imageMem = image->getDeviceMemory(device);
-      std::memcpy(imageSRD, imageMem->cpuSrd(), sizeof(imageSRD));
+    device::Memory* imageMem = image->getDeviceMemory(device);
+    std::memcpy(imageSRD, imageMem->cpuSrd(), sizeof(imageSRD));
 
-      device::Sampler* samplerMem = sampler->getDeviceSampler(device);
-      std::memcpy(samplerSRD, samplerMem->hwState(), sizeof(samplerSRD));
-    }
-  };
+    device::Sampler* samplerMem = sampler->getDeviceSampler(device);
+    std::memcpy(samplerSRD, samplerMem->hwState(), sizeof(samplerSRD));
+  }
 };
 
 amd::Image* ihipImageCreate(const cl_channel_order channelOrder,
@@ -290,11 +288,11 @@ hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject,
   }
 
   void *texObjectBuffer = nullptr;
-  ihipMalloc(&texObjectBuffer, sizeof(hip::TextureObject), CL_MEM_SVM_FINE_GRAIN_BUFFER);
+  ihipMalloc(&texObjectBuffer, sizeof(__hip_texture), CL_MEM_SVM_FINE_GRAIN_BUFFER);
   if (texObjectBuffer == nullptr) {
     return hipErrorOutOfMemory;
   }
-  *pTexObject = reinterpret_cast<hipTextureObject_t>(new (texObjectBuffer) hip::TextureObject{image, sampler, *pResDesc, *pTexDesc, (pResViewDesc != nullptr) ? *pResViewDesc : hipResourceViewDesc{}});
+  *pTexObject = new (texObjectBuffer) __hip_texture{image, sampler, *pResDesc, *pTexDesc, (pResViewDesc != nullptr) ? *pResViewDesc : hipResourceViewDesc{}};
 
   return hipSuccess;
 }
@@ -314,17 +312,16 @@ hipError_t ihipDestroyTextureObject(hipTextureObject_t texObject) {
     return hipErrorInvalidValue;
   }
 
-  hip::TextureObject* hipTexObject = reinterpret_cast<hip::TextureObject*>(texObject);
-  const hipResourceType type = hipTexObject->resDesc.resType;
+  const hipResourceType type = texObject->resDesc.resType;
   const bool isImageFromBuffer = (type == hipResourceTypeLinear) || (type == hipResourceTypePitch2D);
   const bool isImageView = ((type == hipResourceTypeArray) || (type == hipResourceTypeMipmappedArray)) &&
-                           !hipTexObject->image->isParent();
+                           !texObject->image->isParent();
   if (isImageFromBuffer || isImageView) {
-    hipTexObject->image->release();
+    texObject->image->release();
   }
 
   // TODO Should call ihipFree() to not polute the api trace.
-  return hipFree(hipTexObject);
+  return hipFree(texObject);
 }
 
 hipError_t hipDestroyTextureObject(hipTextureObject_t texObject) {
@@ -343,8 +340,7 @@ hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc,
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  hip::TextureObject* hipTexObject = reinterpret_cast<hip::TextureObject*>(texObject);
-  *pResDesc = hipTexObject->resDesc;
+  *pResDesc = texObject->resDesc;
 
   HIP_RETURN(hipSuccess);
 }
@@ -358,8 +354,7 @@ hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  hip::TextureObject* hipTexObject = reinterpret_cast<hip::TextureObject*>(texObject);
-  *pResViewDesc = hipTexObject->resViewDesc;
+  *pResViewDesc = texObject->resViewDesc;
 
   HIP_RETURN(hipSuccess);
 }
@@ -373,8 +368,7 @@ hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc,
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  hip::TextureObject* hipTexObject = reinterpret_cast<hip::TextureObject*>(texObject);
-  *pTexDesc = hipTexObject->texDesc;
+  *pTexDesc = texObject->texDesc;
 
   HIP_RETURN(hipSuccess);
 }