diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h
index 1103bb2274..6459839180 100644
--- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h
+++ b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h
@@ -50,15 +50,64 @@ hipMemcpyHostToHost
 ,hipMemcpyDeviceToDevice
 ,hipMemcpyDefault
 } hipMemcpyKind ;
+#if 0
+typedef enum hipTextureAddressMode {
+    hipAddressModeWrap   = 0,
+    hipAddressModeClamp  = 1,
+    hipAddressModeMirror = 2,
+    hipAddressModeBorder = 3
+}hipTextureAddressMode;
+#endif
+#define hipTextureAddressMode cudaTextureAddressMode
+#define hipAddressModeWrap cudaAddressModeWrap
+#define hipAddressModeClamp cudaAddressModeClamp
+#define hipAddressModeMirror cudaAddressModeMirror
+#define hipAddressModeBorder cudaAddressModeBorder
 
+#if 0
+typedef enum hipTextureFilterMode {
+    hipFilterModePoint  = 0,
+    hipFilterModeLinear = 1
+}hipTextureFilterMode;
+
+#endif
+#if 1
+#define hipTextureFilterMode cudaTextureFilterMode
+#define hipFilterModePoint cudaFilterModePoint
+#define hipFilterModeLinear cudaFilterModeLinear
+#endif
+#if 0
+typedef enum hipTextureReadMode {
+    hipReadModeElementType     = 0,
+    hipReadModeNormalizedFloat = 1
+}hipTextureReadMode;
+#endif
+#define hipTextureReadMode cudaTextureReadMode
+#define hipReadModeElementType cudaReadModeElementType
+#define hipReadModeNormalizedFloat cudaReadModeNormalizedFloat
+
+typedef enum hipChannelFormatKind {
+    hipChannelFormatKindSigned = 0,
+    hipChannelFormatKindUnsigned  = 1,
+    hipChannelFormatKindFloat = 2,
+    hipChannelFormatKindNone = 3
+}hipChannelFormatKind;
+#if 0
+typedef enum hipResourceType {
+    hipResourceTypeArray = 0,
+    hipResourceTypeMipmappedArray = 1,
+    hipResourceTypeLinear = 2,
+    hipResourceTypePitch2D = 3
+}hipResourceType;
+#endif
+#define hipResourceType cudaResourceType
+#define hipResourceTypeArray cudaResourceTypeArray
+#define hipResourceTypeMipmappedArray cudaResourceTypeMipmappedArray
+#define hipResourceTypeLinear cudaResourceTypeLinear
+#define hipResourceTypePitch2D cudaResourceTypePitch2D
+//
 // hipErrorNoDevice.
 
-/*typedef enum hipTextureFilterMode
-{
-    hipFilterModePoint = cudaFilterModePoint,  ///< Point filter mode.
-//! @warning cudaFilterModeLinear is not supported.
-} hipTextureFilterMode;*/
-#define hipFilterModePoint cudaFilterModePoint
 
 //! Flags that can be used with hipEventCreateWithFlags:
 #define hipEventDefault              cudaEventDefault
@@ -119,10 +168,16 @@ typedef CUdevice hipDevice_t;
 typedef CUmodule hipModule_t;
 typedef CUfunction hipFunction_t;
 typedef CUdeviceptr hipDeviceptr_t;
-typedef enum cudaChannelFormatKind hipChannelFormatKind;
-typedef struct cudaChannelFormatDesc hipChannelFormatDesc;
-typedef enum cudaTextureReadMode hipTextureReadMode;
+//typedef enum cudaChannelFormatKind hipChannelFormatKind;
+//typedef struct cudaChannelFormatDesc hipChannelFormatDesc;
 typedef struct cudaArray hipArray;
+typedef struct cudaArray* hipArray_const_t;
+//typedef  cudaArray_const_t  hipArray;
+#define hipArrayDefault cudaArrayDefault
+
+typedef cudaTextureObject_t hipTextureObject_t;
+#define hipTextureType2D cudaTextureType2D;
+#define hipDeviceMapHost cudaDeviceMapHost
 
 // Flags that can be used with hipStreamCreateWithFlags
 #define hipStreamDefault            cudaStreamDefault
@@ -130,7 +185,9 @@ typedef struct cudaArray hipArray;
 
 //typedef cudaChannelFormatDesc hipChannelFormatDesc;
 #define hipChannelFormatDesc cudaChannelFormatDesc
-
+#define hipResourceDesc cudaResourceDesc
+#define hipTextureDesc cudaTextureDesc
+#define hipResourceViewDesc cudaResourceViewDesc
 //adding code for hipmemSharedConfig
 #define hipSharedMemBankSizeDefault cudaSharedMemBankSizeDefault
 #define hipSharedMemBankSizeFourByte cudaSharedMemBankSizeFourByte
@@ -214,6 +271,58 @@ inline static enum cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind ki
 }
 }
 
+inline static cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMode(hipTextureAddressMode kind) {
+    switch(kind) {
+    case hipAddressModeWrap:
+        return cudaAddressModeWrap;
+    case hipAddressModeClamp:
+        return cudaAddressModeClamp;
+    case hipAddressModeMirror:
+        return cudaAddressModeMirror;
+    case hipAddressModeBorder:
+        return cudaAddressModeBorder;
+    default:
+        return cudaAddressModeWrap;
+}
+}
+
+inline static cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode(hipTextureFilterMode kind) {
+    switch(kind) {
+    case hipFilterModePoint:
+        return cudaFilterModePoint;
+    case hipFilterModeLinear:
+        return cudaFilterModeLinear;
+    default:
+        return cudaFilterModePoint;
+}
+}
+
+inline static cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTextureReadMode kind) {
+    switch(kind) {
+    case hipReadModeElementType:
+        return cudaReadModeElementType;
+    case hipReadModeNormalizedFloat:
+        return cudaReadModeNormalizedFloat;
+    default:
+        return cudaReadModeElementType;
+}
+}
+
+inline static cudaChannelFormatKind hipChannelFormatKindToCudaChannelFormatKind(hipChannelFormatKind kind) {
+    switch(kind) {
+    case hipChannelFormatKindSigned :
+        return cudaChannelFormatKindSigned ;
+    case hipChannelFormatKindUnsigned :
+        return cudaChannelFormatKindUnsigned ;
+    case hipChannelFormatKindFloat :
+        return cudaChannelFormatKindFloat  ;
+    case hipChannelFormatKindNone  :
+        return cudaChannelFormatKindNone ;
+    default:
+        return cudaChannelFormatKindNone ;
+}
+}
+
 /**
  * Stream CallBack struct
  */
@@ -262,9 +371,17 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla
 	return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags));
 }
 
-inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) {
-  return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags));
+#if __cplusplus
+inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
+                          size_t width, size_t height = 0, unsigned int flags = hipArrayDefault) {
+    return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags));
 }
+#else
+inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc,
+                          size_t width, size_t height, unsigned int flags) {
+    return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags));
+}
+#endif
 
 inline static hipError_t hipFreeArray(hipArray* array) {
   return hipCUDAErrorTohipError(cudaFreeArray(array));
@@ -1007,11 +1124,44 @@ inline static hipError_t  hipUnbindTexture(struct texture<T, dim, readMode> *tex
 		return  hipCUDAErrorTohipError(cudaUnbindTexture(tex));
 }
 
+template <class T, int dim, enum hipTextureReadMode readMode>
+inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex,
+                                 hipArray_const_t array,
+                                 const struct hipChannelFormatDesc& desc)
+{
+    return  hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array,desc));
+}
+
+template <class T, int dim, enum cudaTextureReadMode readMode>
+inline static  hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> &tex,
+                                 hipArray_const_t array)
+{
+    return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array));
+}
+
 template <class T>
 inline static hipChannelFormatDesc hipCreateChannelDesc()
 {
 		return cudaCreateChannelDesc<T>();
 }
+
+hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f)
+{
+    return cudaCreateChannelDesc(x , y , z , w, hipChannelFormatKindToCudaChannelFormatKind(f)); 
+}
+
+inline static hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
+                                  const hipResourceDesc* pResDesc,
+                                  const hipTextureDesc* pTexDesc,
+                                  const hipResourceViewDesc* pResViewDesc)
+{
+    return hipCUDAErrorTohipError(cudaCreateTextureObject(pTexObject, pResDesc, pTexDesc, pResViewDesc));
+}
+
+inline static hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject)
+{
+    return hipCUDAErrorTohipError(cudaDestroyTextureObject(textureObject));
+}
 #endif //__CUDACC__
 
 #endif //HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H