Headers need to export C symbols for texture API

This also adds declarations of all the missing texture APIs.

hipTexRefSet*() functions need to take a textureReference as a ptr for type erasure to work. Runtime has been modified to accomodate this.

This change only applies to VDI.

Change-Id: Icf43cc5bd44dfc2c39084b7fe56d5a793bf7319f

hipamd/include/hip/hcc_detail/channel_descriptor.h

@@ -29,8 +29,14 @@ THE SOFTWARE.
 
 #ifdef __cplusplus
 
+#if __HIP_VDI__
+extern "C" {
+#endif
 HIP_PUBLIC_API
 hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f);
+#if __HIP_VDI__
+}
+#endif
 
 static inline hipChannelFormatDesc hipCreateChannelDescHalf() {
     int e = (int)sizeof(unsigned short) * 8;

hipamd/include/hip/hcc_detail/hip_runtime_api.h

@@ -3314,6 +3314,172 @@ hipError_t hipLaunchKernel(const void* function_address,
                            size_t sharedMemBytes __dparm(0),
                            hipStream_t stream __dparm(0));
 
+#if __HIP_VDI__
+hipError_t hipBindTexture(
+    size_t* offset,
+    const textureReference* tex,
+    const void* devPtr,
+    const hipChannelFormatDesc* desc,
+    size_t size = UINT_MAX);
+
+hipError_t hipBindTexture2D(
+    size_t* offset,
+    const textureReference* tex,
+    const void* devPtr,
+    const hipChannelFormatDesc* desc,
+    size_t width,
+    size_t height,
+    size_t pitch);
+
+hipError_t hipBindTextureToArray(
+    const textureReference* tex,
+    hipArray_const_t array,
+    const hipChannelFormatDesc* desc);
+
+hipError_t hipBindTextureToMipmappedArray(
+    const textureReference* tex,
+    hipMipmappedArray_const_t mipmappedArray,
+    const hipChannelFormatDesc* desc);
+
+hipError_t hipGetTextureAlignmentOffset(
+    size_t* offset,
+    const textureReference* texref);
+
+hipError_t hipGetTextureReference(
+    const textureReference** texref,
+    const void* symbol);
+
+hipError_t hipUnbindTexture(const textureReference* tex);
+
+hipError_t hipCreateTextureObject(
+    hipTextureObject_t* pTexObject,
+    const hipResourceDesc* pResDesc,
+    const hipTextureDesc* pTexDesc,
+    const hipResourceViewDesc* pResViewDesc);
+
+hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject);
+
+hipError_t hipGetChannelDesc(
+    hipChannelFormatDesc* desc,
+    hipArray_const_t array);
+
+hipError_t hipGetTextureObjectResourceDesc(
+    hipResourceDesc* pResDesc,
+    hipTextureObject_t textureObject);
+
+hipError_t hipGetTextureObjectResourceViewDesc(
+    hipResourceViewDesc* pResViewDesc,
+    hipTextureObject_t textureObject);
+
+hipError_t hipGetTextureObjectTextureDesc(
+    hipTextureDesc* pTexDesc,
+    hipTextureObject_t textureObject);
+
+hipError_t hipTexRefGetAddress(
+    hipDeviceptr_t* dev_ptr,
+    const textureReference* texRef);
+
+hipError_t hipTexRefGetAddressMode(
+    hipTextureAddressMode* pam,
+    const textureReference* texRef,
+    int dim);
+
+hipError_t hipTexRefGetFilterMode(
+    hipTextureFilterMode* pfm,
+    const textureReference* texRef);
+
+hipError_t hipTexRefGetFlags(
+    unsigned int* pFlags,
+    const textureReference* texRef);
+
+hipError_t hipTexRefGetFormat(
+    hipArray_Format* pFormat,
+    int* pNumChannels,
+    const textureReference* texRef);
+
+hipError_t hipTexRefGetMaxAnisotropy(
+    int* pmaxAnsio,
+    const textureReference* texRef);
+
+hipError_t hipTexRefGetMipmapFilterMode(
+    hipTextureFilterMode* pfm,
+    const textureReference* texRef);
+
+hipError_t hipTexRefGetMipmapLevelBias(
+    float* pbias,
+    const textureReference* texRef);
+
+hipError_t hipTexRefGetMipmapLevelClamp(
+    float* pminMipmapLevelClamp,
+    float* pmaxMipmapLevelClamp,
+    const textureReference* texRef);
+
+hipError_t hipTexRefGetMipMappedArray(
+    hipMipmappedArray_t* pArray,
+    const textureReference* texRef);
+
+hipError_t hipTexRefSetAddress(
+    size_t* ByteOffset,
+    textureReference* texRef,
+    hipDeviceptr_t dptr,
+    size_t bytes);
+
+hipError_t hipTexRefSetAddress2D(
+    textureReference* texRef,
+    const HIP_ARRAY_DESCRIPTOR* desc,
+    hipDeviceptr_t dptr,
+    size_t Pitch);
+
+hipError_t hipTexRefSetAddressMode(
+    textureReference* texRef,
+    int dim,
+    hipTextureAddressMode am);
+
+hipError_t hipTexRefSetArray(
+    textureReference* tex,
+    hipArray_const_t array,
+    unsigned int flags);
+
+hipError_t hipTexRefSetBorderColor(
+    textureReference* texRef,
+    float* pBorderColor);
+
+hipError_t hipTexRefSetFilterMode(
+    textureReference* texRef,
+    hipTextureFilterMode fm);
+
+hipError_t hipTexRefSetFlags(
+    textureReference* texRef,
+    unsigned int Flags);
+
+hipError_t hipTexRefSetFormat(
+    textureReference* texRef,
+    hipArray_Format fmt,
+    int NumPackedComponents);
+
+hipError_t hipTexRefSetMaxAnisotropy(
+    textureReference* texRef,
+    unsigned int maxAniso);
+
+hipError_t hipTexRefSetMipmapFilterMode(
+    textureReference* texRef,
+    hipTextureFilterMode fm);
+
+hipError_t hipTexRefSetMipmapLevelBias(
+    textureReference* texRef,
+    float bias);
+
+hipError_t hipTexRefSetMipmapLevelClamp(
+    textureReference* texRef,
+    float minMipMapLevelClamp,
+    float maxMipMapLevelClamp);
+
+hipError_t hipTexRefSetMipmappedArray(
+    textureReference* texRef,
+    hipMipmappedArray* mipmappedArray,
+    unsigned int Flags);
+#endif
+
 /**
  * @}
  */
@@ -3362,14 +3528,7 @@ const char* hipKernelNameRef(const hipFunction_t f);
 
 class TlsData;
 
-#if __HIP_VDI__
-hipError_t hipBindTexture(
-    size_t* offset,
-    const textureReference* tex,
-    const void* devPtr,
-    const hipChannelFormatDesc* desc,
-    size_t size = UINT_MAX);
-#else
+#if !__HIP_VDI__
 hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr,
                           const hipChannelFormatDesc* desc, size_t size = UINT_MAX);
 #endif
@@ -3426,16 +3585,7 @@ hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& tex,
 #endif
 
 // C API
-#if __HIP_VDI__
-hipError_t hipBindTexture2D(
-    size_t* offset,
-    const textureReference* tex,
-    const void* devPtr,
-    const hipChannelFormatDesc* desc,
-    size_t width,
-    size_t height,
-    size_t pitch);
-#else
+#if !__HIP_VDI__
 hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* devPtr,
                             const hipChannelFormatDesc* desc, size_t width, size_t height,
                             size_t pitch);
@@ -3466,12 +3616,7 @@ hipError_t hipBindTexture2D(size_t* offset, struct texture<T, dim, readMode>& te
 #endif
 
 // C API
-#if __HIP_VDI__
-hipError_t hipBindTextureToArray(
-    const textureReference* tex,
-    hipArray_const_t array,
-    const hipChannelFormatDesc* desc);
-#else
+#if !__HIP_VDI__
 hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array,
                                  const hipChannelFormatDesc* desc);
 #endif
@@ -3508,9 +3653,11 @@ inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>
 #endif
 
 // C API
+#if !__HIP_VDI__
 hipError_t hipBindTextureToMipmappedArray(const textureReference* tex,
                                           hipMipmappedArray_const_t mipmappedArray,
                                           const hipChannelFormatDesc* desc);
+#endif
 
 #if !__HIP_VDI__
 template <class T, int dim, enum hipTextureReadMode readMode>
@@ -3566,7 +3713,9 @@ inline hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchPara
  *
  *  @return #hipSuccess
  **/
+#if !__HIP_VDI__
 hipError_t hipUnbindTexture(const textureReference* tex);
+#endif
 
 #if !__HIP_VDI__
 extern hipError_t ihipUnbindTextureImpl(const hipTextureObject_t& textureObject);
@@ -3579,6 +3728,7 @@ hipError_t hipUnbindTexture(struct texture<T, dim, readMode>& tex) {
 }
 #endif
 
+#if !__HIP_VDI__
 hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array);
 hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* texref);
 hipError_t hipGetTextureReference(const textureReference** texref, const void* symbol);
@@ -3616,6 +3766,7 @@ hipError_t hipTexRefGetAddress(hipDeviceptr_t* dev_ptr, textureReference tex);
 
 hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPTOR* desc,
                                  hipDeviceptr_t devPtr, size_t pitch);
+#endif
 
 hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, const hipResourceDesc* pResDesc);
 

hipamd/vdi/hip_hcc.def.in

@@ -173,31 +173,45 @@ hipRegisterActivityCallback
 hipRemoveActivityCallback
 hipApiName
 hipKernelNameRef
-hipCreateTextureObject
-hipDestroyTextureObject
-hipGetTextureObjectResourceDesc
-hipGetTextureObjectResourceViewDesc
-hipGetTextureObjectTextureDesc
 hipBindTexture
 hipBindTexture2D
 hipBindTextureToArray
 hipBindTextureToMipmappedArray
-hipUnbindTexture
-hipGetChannelDesc
 hipGetTextureAlignmentOffset
 hipGetTextureReference
-hipTexRefSetFormat
-hipTexRefSetFlags
-hipTexRefSetFilterMode
-hipTexRefSetAddressMode
-hipTexRefSetArray
-hipTexRefSetAddress
-hipTexRefSetAddress2D
+hipUnbindTexture
+hipCreateChannelDesc
+hipCreateTextureObject
+hipDestroyTextureObject
+hipGetChannelDesc
+hipGetTextureObjectResourceDesc
+hipGetTextureObjectResourceViewDesc
+hipGetTextureObjectTextureDesc
 hipTexRefGetAddress
 hipTexRefGetAddressMode
 hipTexRefGetArray
+hipTexRefGetBorderColor
+hipTexRefGetFilterMode
+hipTexRefGetFlags
+hipTexRefGetFormat
+hipTexRefGetMaxAnisotropy
+hipTexRefGetMipmapFilterMode
+hipTexRefGetMipmapLevelBias
+hipTexRefGetMipmapLevelClamp
+hipTexRefGetMipmappedArray
+hipTexRefSetAddress
+hipTexRefSetAddress2D
+hipTexRefSetAddressMode
 hipTexRefSetArray
-hipCreateChannelDesc
+hipTexRefSetBorderColor
+hipTexRefSetFilterMode
+hipTexRefSetFlags
+hipTexRefSetFormat
+hipTexRefSetMaxAnisotropy
+hipTexRefSetMipmapFilterMode
+hipTexRefSetMipmapLevelBias
+hipTexRefSetMipmapLevelClamp
+hipTexRefSetMipmappedArray
 hipProfilerStart
 hipProfilerStop
 hipHccGetAccelerator

hipamd/vdi/hip_hcc.map.in

@@ -186,34 +186,48 @@ global:
     hiprtcGetErrorString;
     hiprtcAddNameExpression;
     hiprtcVersion;
+    hipBindTexture;
+    hipBindTexture2D;
+    hipBindTextureToArray;
+    hipBindTextureToMipmappedArray;
+    hipGetTextureAlignmentOffset;
+    hipGetTextureReference;
+    hipUnbindTexture;
+    hipCreateChannelDesc;
+    hipCreateTextureObject;
+    hipDestroyTextureObject;
+    hipGetChannelDesc;
+    hipGetTextureObjectResourceDesc;
+    hipGetTextureObjectResourceViewDesc;
+    hipGetTextureObjectTextureDesc;
+    hipTexRefGetAddress;
+    hipTexRefGetAddressMode;
+    hipTexRefGetArray;
+    hipTexRefGetBorderColor;
+    hipTexRefGetFilterMode;
+    hipTexRefGetFlags;
+    hipTexRefGetFormat;
+    hipTexRefGetMaxAnisotropy;
+    hipTexRefGetMipmapFilterMode;
+    hipTexRefGetMipmapLevelBias;
+    hipTexRefGetMipmapLevelClamp;
+    hipTexRefGetMipmappedArray;
+    hipTexRefSetAddress;
+    hipTexRefSetAddress2D;
+    hipTexRefSetAddressMode;
+    hipTexRefSetArray;
+    hipTexRefSetBorderColor;
+    hipTexRefSetFilterMode;
+    hipTexRefSetFlags;
+    hipTexRefSetFormat;
+    hipTexRefSetMaxAnisotropy;
+    hipTexRefSetMipmapFilterMode;
+    hipTexRefSetMipmapLevelBias;
+    hipTexRefSetMipmapLevelClamp;
+    hipTexRefSetMipmappedArray;
     extern "C++" {
     hip_impl::hipLaunchKernelGGLImpl*;
     hip_impl::demangle*;
-    hipCreateTextureObject*;
-    hipDestroyTextureObject*;
-    hipGetTextureObjectResourceDesc*;
-    hipGetTextureObjectResourceViewDesc*;
-    hipGetTextureObjectTextureDesc*;
-    hipBindTexture*;
-    hipBindTexture2D*;
-    hipBindTextureToArray*;
-    hipBindTextureToMipmappedArray*;
-    hipUnbindTexture*;
-    hipGetChannelDesc*;
-    hipGetTextureAlignmentOffset*;
-    hipGetTextureReference*;
-    hipTexRefSetFormat*;
-    hipTexRefSetFlags*;
-    hipTexRefSetFilterMode*;
-    hipTexRefSetAddressMode*;
-    hipTexRefSetArray*;
-    hipTexRefSetAddress*;
-    hipTexRefSetAddress2D*;
-    hipTexRefGetAddress*;
-    hipTexRefGetAddressMode*;
-    hipTexRefGetArray*;
-    hipTexRefSetArray*;
-    hipCreateChannelDesc*;
     hipHccGetAccelerator*;
     hipHccGetAcceleratorView*;
     hipCreateSurfaceObject*;

hipamd/vdi/hip_texture.cpp

@@ -671,12 +671,13 @@ hipError_t hipTexRefSetFilterMode(textureReference* texRef,
 }
 
 hipError_t hipTexRefGetAddressMode(hipTextureAddressMode* pam,
-                                   textureReference texRef,
+                                   const textureReference* texRef,
                                    int dim) {
   // TODO overload operator<<(ostream&, textureReference&).
-  HIP_INIT_API(hipTexRefGetAddressMode, pam, &texRef, dim);
+  HIP_INIT_API(hipTexRefGetAddressMode, pam, texRef, dim);
 
-  if (pam == nullptr) {
+  if ((pam == nullptr) ||
+      (texRef == nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
@@ -685,7 +686,7 @@ hipError_t hipTexRefGetAddressMode(hipTextureAddressMode* pam,
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  *pam = texRef.addressMode[dim];
+  *pam = texRef->addressMode[dim];
 
   HIP_RETURN(hipSuccess);
 }
@@ -709,17 +710,18 @@ hipError_t hipTexRefSetAddressMode(textureReference* texRef,
 }
 
 hipError_t hipTexRefGetArray(hipArray_t* pArray,
-                             textureReference texRef) {
+                             const textureReference* texRef) {
   // TODO overload operator<<(ostream&, textureReference&).
-  HIP_INIT_API(hipTexRefGetArray, pArray, &texRef);
+  HIP_INIT_API(hipTexRefGetArray, pArray, texRef);
 
-  if (pArray == nullptr) {
+  if ((pArray == nullptr) ||
+      (texRef == nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
   hipResourceDesc resDesc = {};
   // TODO use ihipGetTextureObjectResourceDesc() to not pollute the API trace.
-  hipError_t error = hipGetTextureObjectResourceDesc(&resDesc, texRef.textureObject);
+  hipError_t error = hipGetTextureObjectResourceDesc(&resDesc, texRef->textureObject);
   if (error != hipSuccess) {
     return HIP_RETURN(error);
   }
@@ -769,17 +771,18 @@ hipError_t hipTexRefSetArray(textureReference* texRef,
 }
 
 hipError_t hipTexRefGetAddress(hipDeviceptr_t* dptr,
-                               textureReference texRef) {
+                               const textureReference* texRef) {
   // TODO overload operator<<(ostream&, textureReference&).
-  HIP_INIT_API(hipTexRefGetAddress, dptr, &texRef);
+  HIP_INIT_API(hipTexRefGetAddress, dptr, texRef);
 
-  if (dptr == nullptr) {
+  if ((dptr == nullptr) ||
+      (texRef == nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
   hipResourceDesc resDesc = {};
   // TODO use ihipGetTextureObjectResourceDesc() to not pollute the API trace.
-  hipError_t error = hipGetTextureObjectResourceDesc(&resDesc, texRef.textureObject);
+  hipError_t error = hipGetTextureObjectResourceDesc(&resDesc, texRef->textureObject);
   if (error != hipSuccess) {
     return HIP_RETURN(error);
   }
@@ -869,11 +872,12 @@ hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannel
 }
 
 hipError_t hipTexRefGetBorderColor(float* pBorderColor,
-                                   textureReference texRef) {
+                                   const textureReference* texRef) {
   // TODO overload operator<<(ostream&, textureReference&).
-  HIP_INIT_API(hipTexRefGetBorderColor, pBorderColor, &texRef);
+  HIP_INIT_API(hipTexRefGetBorderColor, pBorderColor, texRef);
 
-  if (pBorderColor == nullptr) {
+  if ((pBorderColor == nullptr) ||
+      (texRef == nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
@@ -885,39 +889,41 @@ hipError_t hipTexRefGetBorderColor(float* pBorderColor,
 }
 
 hipError_t hipTexRefGetFilterMode(hipTextureFilterMode* pfm,
-                                  textureReference texRef) {
+                                  const textureReference* texRef) {
   // TODO overload operator<<(ostream&, textureReference&).
-  HIP_INIT_API(hipTexRefGetFilterMode, pfm, &texRef);
+  HIP_INIT_API(hipTexRefGetFilterMode, pfm, texRef);
 
-  if (pfm == nullptr) {
+  if ((pfm == nullptr) ||
+      (texRef == nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  *pfm = texRef.filterMode;
+  *pfm = texRef->filterMode;
 
   HIP_RETURN(hipSuccess);
 }
 
 hipError_t hipTexRefGetFlags(unsigned int* pFlags,
-                             textureReference texRef) {
+                             const textureReference* texRef) {
   // TODO overload operator<<(ostream&, textureReference&).
-  HIP_INIT_API(hipTexRefGetFlags, pFlags, &texRef);
+  HIP_INIT_API(hipTexRefGetFlags, pFlags, texRef);
 
-  if (pFlags == nullptr) {
+  if ((pFlags == nullptr) ||
+      (texRef == nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
   *pFlags = 0;
 
-  if (texRef.readMode == hipReadModeElementType) {
+  if (texRef->readMode == hipReadModeElementType) {
     *pFlags |= HIP_TRSF_READ_AS_INTEGER;
   }
 
-  if (texRef.normalized == 1) {
+  if (texRef->normalized == 1) {
     *pFlags |= HIP_TRSF_NORMALIZED_COORDINATES;
   }
 
-  if (texRef.sRGB == 1) {
+  if (texRef->sRGB == 1) {
     *pFlags |= HIP_TRSF_SRGB;
   }
 
@@ -926,92 +932,98 @@ hipError_t hipTexRefGetFlags(unsigned int* pFlags,
 
 hipError_t hipTexRefGetFormat(hipArray_Format* pFormat,
                               int* pNumChannels,
-                              textureReference texRef) {
+                              const textureReference* texRef) {
   // TODO overload operator<<(ostream&, textureReference&).
-  HIP_INIT_API(hipTexRefGetFormat, pFormat, pNumChannels, &texRef);
+  HIP_INIT_API(hipTexRefGetFormat, pFormat, pNumChannels, texRef);
 
   if ((pFormat == nullptr) ||
-      (pNumChannels == nullptr)) {
+      (pNumChannels == nullptr) ||
+      (texRef == nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  *pFormat = texRef.format;
-  *pNumChannels = texRef.numChannels;
+  *pFormat = texRef->format;
+  *pNumChannels = texRef->numChannels;
 
   HIP_RETURN(hipSuccess);
 }
 
 hipError_t hipTexRefGetMaxAnisotropy(int* pmaxAnsio,
-                                     textureReference texRef) {
+                                     const textureReference* texRef) {
   // TODO overload operator<<(ostream&, textureReference&).
-  HIP_INIT_API(hipTexRefGetMaxAnisotropy, pmaxAnsio, &texRef);
+  HIP_INIT_API(hipTexRefGetMaxAnisotropy, pmaxAnsio, texRef);
 
-  if (pmaxAnsio == nullptr) {
+  if ((pmaxAnsio == nullptr) ||
+      (texRef == nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  *pmaxAnsio = texRef.maxAnisotropy;
+  *pmaxAnsio = texRef->maxAnisotropy;
 
   HIP_RETURN(hipErrorInvalidValue);
 }
 
 hipError_t hipTexRefGetMipmapFilterMode(hipTextureFilterMode* pfm,
-                                        textureReference texRef) {
+                                        const textureReference* texRef) {
   // TODO overload operator<<(ostream&, textureReference&).
-  HIP_INIT_API(hipTexRefGetMipmapFilterMode, pfm, &texRef);
+  HIP_INIT_API(hipTexRefGetMipmapFilterMode, pfm, texRef);
 
-  if (pfm == nullptr) {
+  if ((pfm == nullptr) ||
+      (texRef == nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  *pfm = texRef.mipmapFilterMode;
+  *pfm = texRef->mipmapFilterMode;
 
   HIP_RETURN(hipErrorInvalidValue);
 }
 
 hipError_t hipTexRefGetMipmapLevelBias(float* pbias,
-                                       textureReference texRef) {
+                                       const textureReference* texRef) {
   // TODO overload operator<<(ostream&, textureReference&).
-  HIP_INIT_API(hipTexRefGetMipmapLevelBias, pbias, &texRef);
+  HIP_INIT_API(hipTexRefGetMipmapLevelBias, pbias, texRef);
 
-  if (pbias == nullptr) {
+  if ((pbias == nullptr) ||
+      (texRef == nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  *pbias = texRef.mipmapLevelBias;
+  *pbias = texRef->mipmapLevelBias;
 
   HIP_RETURN(hipErrorInvalidValue);
 }
 
 hipError_t hipTexRefGetMipmapLevelClamp(float* pminMipmapLevelClamp,
                                         float* pmaxMipmapLevelClamp,
-                                        textureReference texRef) {
+                                        const textureReference* texRef) {
   // TODO overload operator<<(ostream&, textureReference&).
-  HIP_INIT_API(hipTexRefGetMipmapLevelClamp, pminMipmapLevelClamp, pmaxMipmapLevelClamp, &texRef);
+  HIP_INIT_API(hipTexRefGetMipmapLevelClamp, pminMipmapLevelClamp, pmaxMipmapLevelClamp, texRef);
 
   if ((pminMipmapLevelClamp == nullptr) ||
-      (pmaxMipmapLevelClamp == nullptr)){
+      (pmaxMipmapLevelClamp == nullptr) ||
+      (texRef == nullptr)){
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  *pminMipmapLevelClamp = texRef.minMipmapLevelClamp;
-  *pmaxMipmapLevelClamp = texRef.maxMipmapLevelClamp;
+  *pminMipmapLevelClamp = texRef->minMipmapLevelClamp;
+  *pmaxMipmapLevelClamp = texRef->maxMipmapLevelClamp;
 
   HIP_RETURN(hipErrorInvalidValue);
 }
 
-hipError_t hipTexRefGetMipMappedArray(hipMipmappedArray_t* pArray,
-                                      textureReference texRef) {
+hipError_t hipTexRefGetMipmappedArray(hipMipmappedArray_t* pArray,
+                                      const textureReference* texRef) {
   // TODO overload operator<<(ostream&, textureReference&).
-  HIP_INIT_API(hipTexRefGetMipMappedArray, pArray, &texRef);
+  HIP_INIT_API(hipTexRefGetMipmappedArray, pArray, &texRef);
 
-  if (pArray == nullptr) {
+  if ((pArray == nullptr) ||
+      (texRef == nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
   hipResourceDesc resDesc = {};
   // TODO use ihipGetTextureObjectResourceDesc() to not pollute the API trace.
-  hipError_t error = hipGetTextureObjectResourceDesc(&resDesc, texRef.textureObject);
+  hipError_t error = hipGetTextureObjectResourceDesc(&resDesc, texRef->textureObject);
   if (error != hipSuccess) {
     return HIP_RETURN(error);
   }