diff --git a/include/hip/hcc_detail/texture_fetch_functions.h b/include/hip/hcc_detail/texture_fetch_functions.h
index 39ac3ecd89..23a850c52c 100644
--- a/include/hip/hcc_detail/texture_fetch_functions.h
+++ b/include/hip/hcc_detail/texture_fetch_functions.h
@@ -35,7 +35,7 @@ THE SOFTWARE.
     unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
 
 template<typename T>
-struct __hip_is_channel_type
+struct __hip_is_tex_channel_type
 {
     static constexpr bool value =
         std::is_same<T, char>::value ||
@@ -50,17 +50,17 @@ struct __hip_is_channel_type
 template<
     typename T,
     unsigned int rank>
-struct __hip_is_channel_type<HIP_vector_type<T, rank>>
+struct __hip_is_tex_channel_type<HIP_vector_type<T, rank>>
 {
     static constexpr bool value =
-        __hip_is_channel_type<T>::value &&
+        __hip_is_tex_channel_type<T>::value &&
         ((rank == 1) ||
          (rank == 2) ||
          (rank == 4));
 };
 
 template<typename T>
-struct __hip_is_normalized_channel_type
+struct __hip_is_tex_normalized_channel_type
 {
     static constexpr bool value =
         std::is_same<T, char>::value ||
@@ -72,10 +72,10 @@ struct __hip_is_normalized_channel_type
 template<
     typename T,
     unsigned int rank>
-struct __hip_is_normalized_channel_type<HIP_vector_type<T, rank>>
+struct __hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>
 {
     static constexpr bool value =
-        __hip_is_normalized_channel_type<T>::value &&
+        __hip_is_tex_normalized_channel_type<T>::value &&
         ((rank == 1) ||
          (rank == 2) ||
          (rank == 4));
@@ -99,7 +99,7 @@ template <typename T>
 struct __hip_tex_ret<
     T,
     hipReadModeElementType,
-    typename std::enable_if<__hip_is_channel_type<T>::value, bool>::type>
+    typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
 {
     using type = T;
 };
@@ -110,7 +110,7 @@ template<
 struct __hip_tex_ret<
     HIP_vector_type<T, rank>,
     hipReadModeElementType,
-    typename std::enable_if<__hip_is_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
+    typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
 {
     using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeElementType>, rank>;
 };
@@ -119,7 +119,7 @@ template<typename T>
 struct __hip_tex_ret<
     T,
     hipReadModeNormalizedFloat,
-    typename std::enable_if<__hip_is_normalized_channel_type<T>::value, bool>::type>
+    typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
 {
     using type = float;
 };
@@ -130,7 +130,7 @@ template<
 struct __hip_tex_ret<
     HIP_vector_type<T, rank>,
     hipReadModeNormalizedFloat,
-    typename std::enable_if<__hip_is_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
+    typename std::enable_if<__hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
 {
     using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeNormalizedFloat>, rank>;
 };
@@ -333,7 +333,7 @@ template <typename T>
 struct __hip_tex2dgather_ret<
     T,
     hipReadModeElementType,
-    typename std::enable_if<__hip_is_channel_type<T>::value, bool>::type>
+    typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
 {
     using type = HIP_vector_type<T, 4>;
 };
@@ -344,7 +344,7 @@ template<
 struct __hip_tex2dgather_ret<
     HIP_vector_type<T, rank>,
     hipReadModeElementType,
-    typename std::enable_if<__hip_is_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
+    typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
 {
     using type = HIP_vector_type<T, 4>;
 };
@@ -353,7 +353,7 @@ template <typename T>
 struct __hip_tex2dgather_ret<
     T,
     hipReadModeNormalizedFloat,
-    typename std::enable_if<__hip_is_normalized_channel_type<T>::value, bool>::type>
+    typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
 {
     using type = float4;
 };
diff --git a/include/hip/hcc_detail/texture_indirect_functions.h b/include/hip/hcc_detail/texture_indirect_functions.h
index 56784f433a..4facce436e 100644
--- a/include/hip/hcc_detail/texture_indirect_functions.h
+++ b/include/hip/hcc_detail/texture_indirect_functions.h
@@ -28,11 +28,40 @@ THE SOFTWARE.
 #include <hip/hip_texture_types.h>
 #include <hip/hcc_detail/ockl_image.h>
 
+#include <type_traits>
+
 #define TEXTURE_OBJECT_PARAMETERS_INIT                                                            \
     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>
+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,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -40,13 +69,17 @@ static __device__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int x)
 {
     *ptr = tex1Dfetch<T>(textureObject, x);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex1D(hipTextureObject_t textureObject, float x)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -54,13 +87,17 @@ static __device__ T tex1D(hipTextureObject_t textureObject, float x)
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x)
 {
     *ptr = tex1D<T>(textureObject, x);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex2D(hipTextureObject_t textureObject, float x, float y)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -68,13 +105,17 @@ static __device__ T tex2D(hipTextureObject_t textureObject, float x, float y)
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex2D(T *ptr, hipTextureObject_t textureObject, float x, float y)
 {
     *ptr = tex2D<T>(textureObject, x, y);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex3D(hipTextureObject_t textureObject, float x, float y, float z)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -82,13 +123,17 @@ static __device__ T tex3D(hipTextureObject_t textureObject, float x, float y, fl
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex3D(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
 {
     *ptr = tex3D<T>(textureObject, x, y, z);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex1DLayered(hipTextureObject_t textureObject, float x, int layer)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -96,13 +141,17 @@ static __device__ T tex1DLayered(hipTextureObject_t textureObject, float x, int
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex1DLayered(T *ptr, hipTextureObject_t textureObject, float x, int layer)
 {
     *ptr = tex1DLayered<T>(textureObject, x, layer);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex2DLayered(hipTextureObject_t textureObject, float x, float y, int layer)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -110,13 +159,17 @@ static __device__ T tex2DLayered(hipTextureObject_t textureObject, float x, floa
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex2DLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer)
 {
     *ptr = tex1DLayered<T>(textureObject, x, y, layer);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__  T texCubemap(hipTextureObject_t textureObject, float x, float y, float z)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -124,13 +177,17 @@ static __device__  T texCubemap(hipTextureObject_t textureObject, float x, float
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void texCubemap(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
 {
     *ptr = texCubemap<T>(textureObject, x, y, z);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T texCubemapLayered(hipTextureObject_t textureObject, float x, float y, float z, int layer)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -138,13 +195,17 @@ static __device__ T texCubemapLayered(hipTextureObject_t textureObject, float x,
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void texCubemapLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer)
 {
     *ptr = texCubemapLayered<T>(textureObject, x, y, z, layer);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex2Dgather(hipTextureObject_t textureObject, float x, float y, int comp = 0)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -173,13 +234,17 @@ static __device__ T tex2Dgather(hipTextureObject_t textureObject, float x, float
     return {};
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex2Dgather(T *ptr, hipTextureObject_t textureObject, float x, float y, int comp = 0)
 {
     *ptr = texCubemapLayered<T>(textureObject, x, y, comp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex1DLod(hipTextureObject_t textureObject, float x, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -187,13 +252,17 @@ static __device__ T tex1DLod(hipTextureObject_t textureObject, float x, float le
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex1DLod(T *ptr, hipTextureObject_t textureObject, float x, float level)
 {
     *ptr = tex1DLod<T>(textureObject, x, level);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex2DLod(hipTextureObject_t textureObject, float x, float y, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -201,13 +270,17 @@ static __device__ T tex2DLod(hipTextureObject_t textureObject, float x, float y,
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex2DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float level)
 {
     *ptr = tex2DLod<T>(textureObject, x, y, level);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex3DLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -215,13 +288,17 @@ static __device__ T tex3DLod(hipTextureObject_t textureObject, float x, float y,
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex3DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
 {
     *ptr = tex3DLod<T>(textureObject, x, y, z, level);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -229,13 +306,17 @@ static __device__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, i
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, int layer, float level)
 {
     *ptr = tex1DLayeredLod<T>(textureObject, x, layer, level);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__  T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -243,13 +324,17 @@ static __device__  T tex2DLayeredLod(hipTextureObject_t textureObject, float x,
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float level)
 {
     *ptr = tex2DLayeredLod<T>(textureObject, x, y, layer, level);
 }
 
-template <class T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T texCubemapLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -257,13 +342,17 @@ static __device__ T texCubemapLod(hipTextureObject_t textureObject, float x, flo
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void texCubemapLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
 {
     *ptr = texCubemapLod<T>(textureObject, x, y, z, level);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -273,13 +362,17 @@ static __device__ T texCubemapGrad(hipTextureObject_t textureObject, float x, fl
     return {};
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void texCubemapGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
 {
     *ptr = texCubemapGrad<T>(textureObject, x, y, z, dPdx, dPdy);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -287,13 +380,17 @@ static __device__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
 {
     *ptr = texCubemapLayeredLod<T>(textureObject, x, y, z, layer, level);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex1DGrad(hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -301,13 +398,17 @@ static __device__ T tex1DGrad(hipTextureObject_t textureObject, float x, float d
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex1DGrad(T *ptr, hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
 {
     *ptr = tex1DGrad<T>(textureObject, x, dPdx, dPdy);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -315,13 +416,17 @@ static __device__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex2DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
 {
     *ptr = tex2DGrad<T>(textureObject, x, y, dPdx, dPdy);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -329,13 +434,17 @@ static __device__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex3DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
 {
     *ptr = tex3DGrad<T>(textureObject, x, y, z, dPdx, dPdy);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -343,13 +452,17 @@ static __device__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x,
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
 {
     *ptr = tex1DLayeredGrad<T>(textureObject, x, layer, dPdx, dPdy);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -357,13 +470,17 @@ static __device__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x,
     return *reinterpret_cast<T*>(&tmp);
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
 {
     *ptr = tex2DLayeredGrad<T>(textureObject, x, y, layer, dPdx, dPdy);
 }
 
-template <class T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__  T texCubemapLayeredGrad(hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -373,7 +490,9 @@ static __device__  T texCubemapLayeredGrad(hipTextureObject_t textureObject, flo
     return {};
 }
 
-template <typename T>
+template <
+    typename T,
+    typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
 static __device__ void texCubemapLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
 {
     *ptr = texCubemapLayeredGrad<T>(textureObject, x, y, z, layer, dPdx, dPdy);