Update hipModuleGetTexRef API

[ROCm/clr commit: 76f174b536]

projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h

@@ -1968,7 +1968,7 @@ hipError_t hipModuleGetFunction(hipFunction_t *function, hipModule_t module, con
  */
 hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name);
 
-hipError_t hipModuleGetTexRef(hipDeviceptr_t *dptr, hipModule_t hmod, const char* name);
+hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name);
 /**
  * @brief builds module from code object which resides in host memory. Image is pointer to that location.
  *

projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp

@@ -71,13 +71,15 @@ bool runTest(int argc, char **argv)
 	copyParam.widthInBytes = copyParam.srcPitch;
 	copyParam.height = height;
     hipMemcpy_2D(&copyParam);
-
-    hipTexRefSetAddressMode(&tex, 0, hipAddressModeWrap);
-	hipTexRefSetAddressMode(&tex, 1, hipAddressModeWrap);
-	hipTexRefSetFilterMode(&tex, hipFilterModePoint);
-    hipTexRefSetFlags(&tex, 0);
-	hipTexRefSetFormat(&tex, HIP_AD_FORMAT_FLOAT, 1);
-    hipTexRefSetArray(&tex, array, HIP_TRSA_OVERRIDE_FORMAT);
+    
+    textureReference* texref;
+    hipModuleGetTexRef(&texref, Module, "tex");
+    hipTexRefSetAddressMode(texref, 0, hipAddressModeWrap);
+	hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap);
+	hipTexRefSetFilterMode(texref, hipFilterModePoint);
+    hipTexRefSetFlags(texref, 0);
+	hipTexRefSetFormat(texref, HIP_AD_FORMAT_FLOAT, 1);
+    hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT);
 
     float* dData = NULL;
     hipMalloc((void **) &dData, size);

projects/clr/hipamd/src/hip_module.cpp

@@ -838,11 +838,11 @@ hipError_t hipModuleLoadDataEx(hipModule_t *module, const void *image, unsigned
     return hipModuleLoadData(module, image);
 }
 
-hipError_t hipModuleGetTexRef(hipDeviceptr_t *dptr, hipModule_t hmod, const char* name)
+hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name)
 {
-    HIP_INIT_API(dptr, hmod, name);
+    HIP_INIT_API(texRef, hmod, name);
     hipError_t ret = hipSuccess;
-    if(dptr == NULL){
+    if(texRef == NULL){
         return ihipLogStatus(hipErrorInvalidValue);
     }
     if(name == NULL || hmod == NULL){
@@ -851,7 +851,7 @@ hipError_t hipModuleGetTexRef(hipDeviceptr_t *dptr, hipModule_t hmod, const char
     else{
         const auto it = coGlobals.find(name);
 		if (it == coGlobals.end()) return ihipLogStatus(hipErrorInvalidValue);
-		*dptr = reinterpret_cast<void*>(it->second);
+		*texRef = reinterpret_cast<textureReference*>(it->second);
 		return ihipLogStatus(ret);
     }
 }