diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h
index 9451ac65eb..d145c7fad5 100644
--- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h
+++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h
@@ -1057,6 +1057,28 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         return HIP_vector_type<T, n>{x} <<= y;
     }
 
+    /*
+     * Map HIP_vector_type<U, rankU> to HIP_vector_type<T, rankT>
+     */
+    template<
+        typename T,
+        unsigned int rankT,
+        typename U,
+        unsigned int rankU>
+    __forceinline__
+    __HOST_DEVICE__
+    typename std::enable_if<
+        (rankT >= 1 && rankT <= 4 && rankU >= 1 && rankU <= 4),
+        const HIP_vector_type<T, rankT>>::type
+    __hipMapVector(const HIP_vector_type<U, rankU> &u) {
+      HIP_vector_type<T, rankT> t; // Initialized to 0
+      if constexpr (rankT >= 1 && rankU >= 1) t.x = static_cast<T>(u.x);
+      if constexpr (rankT >= 2 && rankU >= 2) t.y = static_cast<T>(u.y);
+      if constexpr (rankT >= 3 && rankU >= 3) t.z = static_cast<T>(u.z);
+      if constexpr (rankT >= 4 && rankU >= 4) t.w = static_cast<T>(u.w);
+      return t;
+    };
+
     #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \
         using CUDA_name##1 = HIP_vector_type<T, 1>;\
         using CUDA_name##2 = HIP_vector_type<T, 2>;\
@@ -2030,223 +2052,6 @@ typedef union {
     type r{x, y, z, w};                                                                            \
     return r;                                                                                      \
   }
-
-template<typename T, typename U>
-__HOST_DEVICE__
-__forceinline__
-typename std::enable_if<
-  sizeof(T) / sizeof(typename T::value_type) == 1 &&
-  sizeof(U) / sizeof(typename U::value_type) >= 1, T>::type
-mapElem(const U &u) {
-  return {
-    static_cast<typename T::value_type>(u.x)
-  };
-}
-
-template<typename T, typename U>
-__HOST_DEVICE__
-__forceinline__
-typename std::enable_if<
-  sizeof(T) / sizeof(typename T::value_type) == 2 &&
-  sizeof(U) / sizeof(typename U::value_type) >= 2, T>::type
-mapElem(const U &u) {
-  return {
-    static_cast<typename T::value_type>(u.x),
-    static_cast<typename T::value_type>(u.y)
-  };
-}
-
-template<typename T, typename U>
-__HOST_DEVICE__
-__forceinline__
-typename std::enable_if<
-  sizeof(T) / sizeof(typename T::value_type) == 3 &&
-  sizeof(U) / sizeof(typename U::value_type) >= 3, T>::type
-mapElem(const U &u) {
-  return {
-   static_cast<typename T::value_type>(u.x),
-   static_cast<typename T::value_type>(u.y),
-   static_cast<typename T::value_type>(u.z)
-  };
-}
-
-template<typename T, typename U>
-__HOST_DEVICE__
-__forceinline__
-typename std::enable_if<
-  sizeof(T) / sizeof(typename T::value_type) == 4 &&
-  sizeof(U) / sizeof(typename U::value_type) == 1, T>::type
-mapElem(const U &u) {
-  return {
-    static_cast<typename T::value_type>(u.x),
-    static_cast<typename T::value_type>(0),
-    static_cast<typename T::value_type>(0),
-    static_cast<typename T::value_type>(0)
-  };
-}
-
-template<typename T, typename U>
-__HOST_DEVICE__
-__forceinline__
-typename std::enable_if<
-  sizeof(T) / sizeof(typename T::value_type) == 4 &&
-  sizeof(U) / sizeof(typename U::value_type) == 2, T>::type
-mapElem(const U &u) {
-  return {
-    static_cast<typename T::value_type>(u.x),
-    static_cast<typename T::value_type>(u.y),
-    static_cast<typename T::value_type>(0),
-    static_cast<typename T::value_type>(0)
-  };
-}
-
-template<typename T, typename U>
-__HOST_DEVICE__
-__forceinline__
-typename std::enable_if<
-  sizeof(T) / sizeof(typename T::value_type) == 4 &&
-  sizeof(U) / sizeof(typename U::value_type) == 3, T>::type
-mapElem(const U &u) {
-  return {
-    static_cast<typename T::value_type>(u.x),
-    static_cast<typename T::value_type>(u.y),
-    static_cast<typename T::value_type>(u.z),
-    static_cast<typename T::value_type>(0)
-  };
-}
-
-template<typename T, typename U>
-__HOST_DEVICE__
-__forceinline__
-typename std::enable_if<
-  sizeof(T) / sizeof(typename T::value_type) == 4 &&
-  sizeof(U) / sizeof(typename U::value_type) >= 4, T>::type
-mapElem(const U &u) {
-  return {
-    static_cast<typename T::value_type>(u.x),
-    static_cast<typename T::value_type>(u.y),
-    static_cast<typename T::value_type>(u.z),
-    static_cast<typename T::value_type>(u.w)
-  };
-}
-
-template<typename T, typename U>
-__HOST_DEVICE__
-__forceinline__
-typename std::enable_if<
-        std::is_same<T, char>::value ||
-        std::is_same<T, unsigned char>::value ||
-        std::is_same<T, short>::value ||
-        std::is_same<T, unsigned short>::value ||
-        std::is_same<T, int>::value ||
-        std::is_same<T, unsigned int>::value ||
-        std::is_same<T, float>::value, const T>::type
-mapFrom(const U &u) {
-  union {
-    U u;
-    T t;
-  } d = { u };
-  return d.t;
-}
-
-template<typename T, typename U>
-__HOST_DEVICE__
-__forceinline__
-typename std::enable_if<
-        (sizeof(T) == sizeof(typename T::value_type)) ||
-        std::is_same<typename T::value_type, int>::value ||
-        std::is_same<typename T::value_type, unsigned int>::value ||
-        std::is_same<typename T::value_type, float>::value, const T>::type
-mapFrom(const U &u) {
-  union {
-    U u;
-    T t;
-  } d = { u };
-  return d.t;
-}
-
-template<typename T, typename U>
-__HOST_DEVICE__
-__forceinline__
-typename std::enable_if<
-        (sizeof(T) > sizeof(typename T::value_type)) && (
-        std::is_same<typename T::value_type, char>::value ||
-        std::is_same<typename T::value_type, unsigned char>::value ||
-        std::is_same<typename T::value_type, short>::value ||
-        std::is_same<typename T::value_type, unsigned short>::value), const T>::type
-mapFrom(const U &u) {
-  union {
-    U u;
-    int4 i4;
-    uint4 u4;
-  } d = { u };
-  if (std::is_signed<typename T::value_type>::value) {
-    return mapElem<T>(d.i4) ;
-  } else {
-    return mapElem<T>(d.u4);
-  }
-}
-
-template<typename U, typename T>
-__HOST_DEVICE__
-__forceinline__
-typename std::enable_if<
-        std::is_same<T, char>::value ||
-        std::is_same<T, unsigned char>::value ||
-        std::is_same<T, short>::value ||
-        std::is_same<T, unsigned short>::value ||
-        std::is_same<T, int>::value ||
-        std::is_same<T, unsigned int>::value ||
-        std::is_same<T, float>::value, const U>::type
-mapTo(const T &t) {
-  union {
-    U u;
-    T t;
-  } d = { 0 };
-  d.t = t;
-  return d.u;
-}
-
-template<typename U, typename T>
-__HOST_DEVICE__
-__forceinline__
-typename std::enable_if<
-        (sizeof(T) == sizeof(typename T::value_type)) ||
-        std::is_same<typename T::value_type, int>::value ||
-        std::is_same<typename T::value_type, unsigned int>::value ||
-        std::is_same<typename T::value_type, float>::value, const U>::type
-mapTo(const T &t) {
-  union {
-    U u;
-    T t;
-  } d = { 0 };
-  d.t = t;
-  return d.u;
-}
-
-template<typename U, typename T>
-__HOST_DEVICE__
-__forceinline__
-typename std::enable_if<
-        (sizeof(T) > sizeof(typename T::value_type)) && (
-        std::is_same<typename T::value_type, char>::value ||
-        std::is_same<typename T::value_type, unsigned char>::value ||
-        std::is_same<typename T::value_type, short>::value ||
-        std::is_same<typename T::value_type, unsigned short>::value), const U>::type
-mapTo(const T &t) {
-  union {
-    U u;
-    int4 i4;
-    uint4 u4;
-  } d = { 0 };
-  if (std::is_signed<typename T::value_type>::value) {
-    d.i4 = mapElem<int4>(t);
-  } else {
-    d.u4 = mapElem<uint4>(t);
-  }
-  return d.u;
-}
-
 #else
 #define DECLOP_MAKE_ONE_COMPONENT(comp, type)                                                      \
   static inline __HOST_DEVICE__ type make_##type(comp x) {                                         \
diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h
index c59995ca39..544a249dca 100644
--- a/projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h
+++ b/projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h
@@ -27,36 +27,14 @@ THE SOFTWARE.
 
 #include <hip/surface_types.h>
 #include <hip/hip_vector_types.h>
+#include <hip/amd_detail/texture_fetch_functions.h>
 #include <hip/amd_detail/ockl_image.h>
 
 #define __HIP_SURFACE_OBJECT_PARAMETERS_INIT                                                            \
     unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)surfObj; 
 
-template<typename T>
-struct __hip_is_isurf_channel_type
-{
-    static constexpr bool value =
-        std::is_same<T, char>::value ||
-        std::is_same<T, unsigned char>::value ||
-        std::is_same<T, short>::value ||
-        std::is_same<T, unsigned short>::value ||
-        std::is_same<T, int>::value ||
-        std::is_same<T, unsigned int>::value ||
-        std::is_same<T, float>::value;
-};
-
-template<
-    typename T,
-    unsigned int rank>
-struct __hip_is_isurf_channel_type<HIP_vector_type<T, rank>>
-{
-    static constexpr bool value =
-        __hip_is_isurf_channel_type<T>::value &&
-        ((rank == 1) ||
-         (rank == 2) ||
-         (rank == 3) ||
-         (rank == 4));
-};
+template <typename T>
+using __hip_is_surf_channel_type = __hip_is_tex_channel_type<T>;
 
 // CUDA is using byte address, need map to pixel address for HIP
 static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format, int order) {
@@ -115,144 +93,144 @@ static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format,
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t surfObj, int x,
         int boundaryMode = hipBoundaryModeZero) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
     auto tmp = __ockl_image_load_1D(i, x);
-    *data = mapFrom<T>(tmp);
+    *data = __hipMapFrom<T>(tmp);
 }
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
-    auto tmp = mapTo<float4::Native_vec_>(data);
+    auto tmp = __hipMapTo<float4::Native_vec_>(data);
     __ockl_image_store_1D(i, x, tmp);
 }
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
     auto tmp = __ockl_image_load_2D(i, int2(x, y).data);
-    *data = mapFrom<T>(tmp);
+    *data = __hipMapFrom<T>(tmp);
 }
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
-    auto tmp = mapTo<float4::Native_vec_>(data);
+    auto tmp = __hipMapTo<float4::Native_vec_>(data);
     __ockl_image_store_2D(i, int2(x, y).data, tmp);
 }
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
     auto tmp = __ockl_image_load_3D(i, int4(x, y, z, 0).data);
-    *data = mapFrom<T>(tmp);
+    *data = __hipMapFrom<T>(tmp);
 }
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
-    auto tmp = mapTo<float4::Native_vec_>(data);
+    auto tmp = __hipMapTo<float4::Native_vec_>(data);
     __ockl_image_store_3D(i, int4(x, y, z, 0).data, tmp);
 }
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
     auto tmp = __ockl_image_load_lod_1D(i, x, layer);
-    *data = mapFrom<T>(tmp);
+    *data = __hipMapFrom<T>(tmp);
 }
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
-    auto tmp = mapTo<float4::Native_vec_>(data);
+    auto tmp = __hipMapTo<float4::Native_vec_>(data);
     __ockl_image_store_lod_1D(i, x, layer, tmp);
 }
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
     auto tmp = __ockl_image_load_lod_2D(i, int2(x, y).data, layer);
-    *data = mapFrom<T>(tmp);
+    *data = __hipMapFrom<T>(tmp);
 }
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
-    auto tmp = mapTo<float4::Native_vec_>(data);
+    auto tmp = __hipMapTo<float4::Native_vec_>(data);
     __ockl_image_store_lod_2D(i, int2(x, y).data, layer, tmp);
 }
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
     auto tmp = __ockl_image_load_CM(i, int2(x, y).data, face);
-    *data = mapFrom<T>(tmp);
+    *data = __hipMapFrom<T>(tmp);
 }
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
-    auto tmp = mapTo<float4::Native_vec_>(data);
+    auto tmp = __hipMapTo<float4::Native_vec_>(data);
     __ockl_image_store_CM(i, int2(x, y).data, face, tmp);
 }
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
         int layer) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
     auto tmp = __ockl_image_load_lod_CM(i, int2(x, y).data, face, layer);
-    *data = mapFrom<T>(tmp);
+    *data = __hipMapFrom<T>(tmp);
 }
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
+    typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
         int layer) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
-    auto tmp = mapTo<float4::Native_vec_>(data);
+    auto tmp = __hipMapTo<float4::Native_vec_>(data);
     __ockl_image_store_lod_CM(i, int2(x, y).data, face, layer, tmp);
 }
 
diff --git a/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h b/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h
index dcf5f2fdf8..a252c0efde 100644
--- a/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h
+++ b/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h
@@ -92,6 +92,100 @@ struct __hip_tex_ret
     static_assert(std::is_same<Enable, void>::value, "Invalid channel type!");
 };
 
+/*
+ * Map from device function return U to scalar texture type T
+ */
+template<typename T, typename U>
+__forceinline__ __device__
+typename std::enable_if<
+  __hip_is_tex_channel_type<T>::value && std::is_scalar<T>::value, const T>::type
+__hipMapFrom(const U &u) {
+  if constexpr (sizeof(T) < sizeof(float)) {
+    union {
+      U u;
+      int i;
+    } d = { u };
+    return static_cast<T>(d.i);
+  } else { // sizeof(T) == sizeof(float)
+    union {
+      U u;
+      T t;
+    } d = { u };
+    return d.t;
+  }
+}
+
+/*
+ * Map from device function return U to vector texture type T
+ */
+template<typename T, typename U>
+__forceinline__ __device__
+typename std::enable_if<__hip_is_tex_channel_type<typename T::value_type>::value, const T>::type
+__hipMapFrom(const U &u) {
+  if constexpr (sizeof(typename T::value_type) < sizeof(float)) {
+    union {
+      U u;
+      int4 i4;
+    } d = { u };
+    return __hipMapVector<typename T::value_type, sizeof(T)/sizeof(typename T::value_type)>(d.i4);
+  } else { // sizeof(typename T::value_type) == sizeof(float)
+    union {
+      U u;
+      T t;
+    } d = { u };
+    return d.t;
+  }
+}
+
+/*
+ * Map from scalar texture type T to device function input U
+ */
+template<typename U, typename T>
+__forceinline__ __device__
+typename std::enable_if<
+  __hip_is_tex_channel_type<T>::value  && std::is_scalar<T>::value, const U>::type
+__hipMapTo(const T &t) {
+  if constexpr (sizeof(T) < sizeof(float)) {
+    union {
+      U u;
+      int i;
+    } d = { 0 };
+    d.i = static_cast<int>(t);
+    return d.u;
+  } else { // sizeof(T) == sizeof(float)
+    union {
+      U u;
+      T t;
+    } d = { 0 };
+    d.t = t;
+    return d.u;
+  }
+}
+
+/*
+ * Map from vector texture type T to device function input U
+ */
+template<typename U, typename T>
+__forceinline__ __device__
+typename std::enable_if<__hip_is_tex_channel_type<typename T::value_type>::value, const U>::type
+__hipMapTo(const T &t) {
+  if constexpr (sizeof(typename T::value_type) < sizeof(float)) {
+    union {
+      U u;
+      int4 i4;
+    } d = { 0 };
+    d.i4 = __hipMapVector<int, 4>(t);
+    return d.u;
+  } else { // sizeof(typename T::value_type) == sizeof(float)
+    union {
+      U u;
+      T t;
+    } d = { 0 };
+    d.t = t;
+    return d.u;
+  }
+}
+
 template <
     typename T,
     hipTextureReadMode readMode>
@@ -137,12 +231,13 @@ struct __hip_tex_ret<
     using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeNormalizedFloat>, rank>;
 };
 
+
 template <typename T, hipTextureReadMode readMode>
 static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1Dfetch(texture<T, hipTextureType1D, readMode> t, int x)
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_load_1Db(i, x);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -150,7 +245,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_1D(i, s, x);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -158,7 +253,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_2D(i, s, float2(x, y).data);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -166,7 +261,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -174,7 +269,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_2Da(i, s, float4(x, y, layer, 0.0f).data);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -182,7 +277,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_3D(i, s, float4(x, y, z, 0.0f).data);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -190,7 +285,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_CM(i, s, float4(x, y, z, 0.0f).data);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -198,7 +293,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -206,7 +301,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_lod_2D(i, s, float2(x, y).data, level);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -214,7 +309,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_lod_1Da(i, s, float2(x, layer).data, level);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -222,7 +317,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_lod_2Da(i, s, float4(x, y, layer, 0.0f).data, level);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -230,7 +325,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_lod_3D(i, s, float4(x, y, z, 0.0f).data, level);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -238,7 +333,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_lod_CM(i, s, float4(x, y, z, 0.0f).data, level);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -246,7 +341,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_CMa(i, s, float4(x, y, z, layer).data);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -254,7 +349,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_lod_CMa(i, s, float4(x, y, z, layer).data, level);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -263,7 +358,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
     TEXTURE_PARAMETERS_INIT;
     // TODO missing in device libs.
     // auto tmp = __ockl_image_sample_grad_CM(i, s, float4(x, y, z, 0.0f).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
-    // return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    // return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
     return {};
 }
 
@@ -273,7 +368,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
     TEXTURE_PARAMETERS_INIT;
     // TODO missing in device libs.
     // auto tmp = __ockl_image_sample_grad_CMa(i, s, float4(x, y, z, layer).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
-    // return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    // return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
     return {};
 }
 
@@ -282,7 +377,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -290,7 +385,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_grad_2D(i, s, float2(x, y).data, float2(dPdx.x, dPdx.y).data,  float2(dPdy.x, dPdy.y).data);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -298,7 +393,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_grad_1Da(i, s, float2(x, layer).data, dPdx, dPdy);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -306,7 +401,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_grad_2Da(i, s, float4(x, y, layer, 0.0f).data, float2(dPdx.x, dPdx.y).data, float2(dPdy.x, dPdy.y).data);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>
@@ -314,7 +409,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_sample_grad_3D(i, s, float4(x, y, z, 0.0f).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
-    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
+    return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <
@@ -367,19 +462,19 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex2dgather_ret_t<T, rea
     switch (comp) {
     case 1: {
         auto tmp = __ockl_image_gather4g_2D(i, s, float2(x, y).data);
-        return mapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
+        return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
     }
     case 2: {
         auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data);
-        return mapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
+        return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
     }
     case 3: {
         auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data);
-        return mapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
+        return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
     }
     default: {
         auto tmp = __ockl_image_gather4r_2D(i, s, float2(x, y).data);
-        return mapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
+        return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
     }
     }
     return {};
diff --git a/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h b/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h
index 0dd04d74dd..af0bfa4cec 100644
--- a/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h
+++ b/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h
@@ -26,6 +26,7 @@ THE SOFTWARE.
 
 #include <hip/hip_vector_types.h>
 #include <hip/hip_texture_types.h>
+#include <hip/amd_detail/texture_fetch_functions.h>
 #include <hip/amd_detail/ockl_image.h>
 
 #if !defined(__HIPCC_RTC__)
@@ -36,30 +37,8 @@ THE SOFTWARE.
     unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \
     unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
 
-template<typename T>
-struct __hip_is_itex_channel_type
-{
-    static constexpr bool value =
-        std::is_same<T, char>::value ||
-        std::is_same<T, unsigned char>::value ||
-        std::is_same<T, short>::value ||
-        std::is_same<T, unsigned short>::value ||
-        std::is_same<T, int>::value ||
-        std::is_same<T, unsigned int>::value ||
-        std::is_same<T, float>::value;
-};
-
-template<
-    typename T,
-    unsigned int rank>
-struct __hip_is_itex_channel_type<HIP_vector_type<T, rank>>
-{
-    static constexpr bool value =
-        __hip_is_itex_channel_type<T>::value &&
-        ((rank == 1) ||
-         (rank == 2) ||
-         (rank == 4));
-};
+template <typename T>
+using __hip_is_itex_channel_type = __hip_is_tex_channel_type<T>;
 
 template <
     typename T,
@@ -68,7 +47,7 @@ static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject,
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_load_1Db(i, x);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -86,7 +65,7 @@ static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject, floa
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_1D(i, s, x);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -104,7 +83,7 @@ static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, floa
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_2D(i, s, float2(x, y).data);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -122,7 +101,7 @@ static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, floa
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_3D(i, s, float4(x, y, z, 0.0f).data);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -140,7 +119,7 @@ static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObjec
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -158,7 +137,7 @@ static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObjec
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_2Da(i, s, float4(x, y, layer, 0.0f).data);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -176,7 +155,7 @@ static __device__ __hip_img_chk__  T texCubemap(hipTextureObject_t textureObject
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_CM(i, s, float4(x, y, z, 0.0f).data);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -194,7 +173,7 @@ static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t texture
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_CMa(i, s, float4(x, y, z, layer).data);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -214,22 +193,22 @@ static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject
     switch (comp) {
     case 1: {
         auto tmp = __ockl_image_gather4r_2D(i, s, float2(x, y).data);
-        return mapFrom<T>(tmp);
+        return __hipMapFrom<T>(tmp);
         break;
     }
     case 2: {
         auto tmp = __ockl_image_gather4g_2D(i, s, float2(x, y).data);
-        return mapFrom<T>(tmp);
+        return __hipMapFrom<T>(tmp);
         break;
     }
     case 3: {
         auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data);
-        return mapFrom<T>(tmp);
+        return __hipMapFrom<T>(tmp);
         break;
     }
     default: {
         auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data);
-        return mapFrom<T>(tmp);
+        return __hipMapFrom<T>(tmp);
         break;
     }
     };
@@ -251,7 +230,7 @@ static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, f
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -269,7 +248,7 @@ static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, f
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_lod_2D(i, s, float2(x, y).data, level);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -287,7 +266,7 @@ static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject, f
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_lod_3D(i, s, float4(x, y, z, 0.0f).data, level);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -305,7 +284,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureOb
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -323,7 +302,7 @@ static __device__ __hip_img_chk__  T tex2DLayeredLod(hipTextureObject_t textureO
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_2Da(i, s, float4(x, y, layer, 0.0f).data);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -341,7 +320,7 @@ static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObje
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_lod_CM(i, s, float4(x, y, z, 0.0f).data, level);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -360,7 +339,7 @@ static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObj
     TEXTURE_OBJECT_PARAMETERS_INIT
     // TODO missing in device libs.
     // auto tmp = __ockl_image_sample_grad_CM(i, s, float4(x, y, z, 0.0f).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
-    // return mapFrom<T>(tmp);
+    // return __hipMapFrom<T>(tmp);
     return {};
 }
 
@@ -379,7 +358,7 @@ static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t text
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_lod_CMa(i, s, float4(x, y, z, layer).data, level);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -397,7 +376,7 @@ static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject,
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -415,7 +394,7 @@ static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject,
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_grad_2D(i, s, float2(x, y).data, float2(dPdx.x, dPdx.y).data,  float2(dPdy.x, dPdy.y).data);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -433,7 +412,7 @@ static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject,
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_grad_3D(i, s, float4(x, y, z, 0.0f).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -451,7 +430,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureO
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_grad_1Da(i, s, float2(x, layer).data, dPdx, dPdy);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -469,7 +448,7 @@ static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureO
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_sample_grad_2Da(i, s, float4(x, y, layer, 0.0f).data, float2(dPdx.x, dPdx.y).data, float2(dPdy.x, dPdy.y).data);
-    return mapFrom<T>(tmp);
+    return __hipMapFrom<T>(tmp);
 }
 
 template <
@@ -488,7 +467,7 @@ static __device__ __hip_img_chk__  T texCubemapLayeredGrad(hipTextureObject_t te
     TEXTURE_OBJECT_PARAMETERS_INIT
     // TODO missing in device libs.
     // auto tmp = __ockl_image_sample_grad_CMa(i, s, float4(x, y, z, layer).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
-    // return mapFrom<T>(tmp);
+    // return __hipMapFrom<T>(tmp);
     return {};
 }