SWDEV-298434 - Delete texture references resources before deleting the ptr.

Change-Id: Ib616aadd45eced1b10891c908e986bbeddc08288

hipamd/src/hip_global.cpp

@@ -41,6 +41,7 @@ DeviceVar::~DeviceVar() {
 
   if (shadowVptr != nullptr) {
     textureReference* texRef = reinterpret_cast<textureReference*>(shadowVptr);
+    ihipUnbindTexture(texRef);
     delete texRef;
     shadowVptr = nullptr;
   }

hipamd/src/hip_internal.hpp

@@ -371,6 +371,7 @@ 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 void getStreamPerThread(hipStream_t& stream);
+extern hipError_t ihipUnbindTexture(textureReference* texRef);
 
 constexpr bool kOptionChangeable = true;
 constexpr bool kNewDevProg = false;

hipamd/src/hip_texture.cpp

@@ -309,7 +309,6 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
   HIP_RETURN(ihipCreateTextureObject(pTexObject, pResDesc, pTexDesc, pResViewDesc));
 }
 
-
 hipError_t ihipDestroyTextureObject(hipTextureObject_t texObject) {
   if (texObject == nullptr) {
     return hipSuccess;
@@ -332,6 +331,28 @@ hipError_t ihipDestroyTextureObject(hipTextureObject_t texObject) {
   return hipFree(texObject);
 }
 
+hipError_t ihipUnbindTexture(textureReference* texRef) {
+
+  hipError_t hip_error = hipSuccess;
+
+  do {
+    if (texRef == nullptr) {
+      hip_error = hipErrorInvalidValue;
+      break;
+    }
+
+    hip_error = ihipDestroyTextureObject(texRef->textureObject);
+    if (hip_error != hipSuccess) {
+      break;
+    }
+
+    const_cast<textureReference*>(texRef)->textureObject = nullptr;
+
+  } while (0);
+
+  return hip_error;
+}
+
 hipError_t hipDestroyTextureObject(hipTextureObject_t texObject) {
   HIP_INIT_API(hipDestroyTextureObject, texObject);
 
@@ -593,14 +614,7 @@ hipError_t hipBindTextureToMipmappedArray(const textureReference* texref,
 hipError_t hipUnbindTexture(const textureReference* texref) {
   HIP_INIT_API(hipUnbindTexture, texref);
 
-  if (texref == nullptr) {
-    HIP_RETURN(hipErrorInvalidValue);
-  }
-
-  const hipTextureObject_t textureObject = texref->textureObject;
-  const_cast<textureReference*>(texref)->textureObject = nullptr;
-
-  HIP_RETURN(ihipDestroyTextureObject(textureObject));
+  HIP_RETURN(ihipUnbindTexture(const_cast<textureReference*>(texref)));
 }
 
 hipError_t hipBindTexture(size_t* offset,