Add constraints to texture fetch functions

When sampling a pixel the hw always returns a float4. The type in the texture reference controls the bitcast that we perform before returning the sampled pixel. Creating a texture with an unsupported will lead to potential UB.

This change makes it so that it's only possible to use textures with a type that makes sense. Using something like texture<int, hipTextureType1D, hipReadModeNormalizedFloat> will now lead to a compilation error with a message "Invalid channel type!".

Change-Id: I7fde44cb1d4b9737e0c48c28cb59c018c59ccaa2

hipamd/include/hip/hcc_detail/texture_fetch_functions.h

@@ -34,20 +34,106 @@ THE SOFTWARE.
     unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)t.textureObject; \
     unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
 
-template <typename T, hipTextureReadMode readMode>
-struct __hip_tex_ret {};
+template<typename T>
+struct __hip_is_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, hipTextureReadMode readMode>
-using __hip_tex_ret_t = typename __hip_tex_ret<T, readMode>::type;
-
-template <typename T>
-struct __hip_tex_ret<T, hipReadModeElementType> { using type = T; };
+template<
+    typename T,
+    unsigned int rank>
+struct __hip_is_channel_type<HIP_vector_type<T, rank>>
+{
+    static constexpr bool value =
+        __hip_is_channel_type<T>::value &&
+        ((rank == 1) ||
+         (rank == 2) ||
+         (rank == 4));
+};
 
 template<typename T>
-struct __hip_tex_ret<T, hipReadModeNormalizedFloat> { using type = float; };
+struct __hip_is_normalized_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;
+};
 
-template<typename T, unsigned int rank>
-struct __hip_tex_ret<HIP_vector_type<T, rank>, hipReadModeNormalizedFloat> { using type = HIP_vector_type<float, rank>; };
+template<
+    typename T,
+    unsigned int rank>
+struct __hip_is_normalized_channel_type<HIP_vector_type<T, rank>>
+{
+    static constexpr bool value =
+        __hip_is_normalized_channel_type<T>::value &&
+        ((rank == 1) ||
+         (rank == 2) ||
+         (rank == 4));
+};
+
+template <
+    typename T,
+    hipTextureReadMode readMode,
+    typename Enable = void>
+struct __hip_tex_ret
+{
+    static_assert(std::is_same<Enable, void>::value, "Invalid channel type!");
+};
+
+template <
+    typename T,
+    hipTextureReadMode readMode>
+using __hip_tex_ret_t = typename __hip_tex_ret<T, readMode, bool>::type;
+
+template <typename T>
+struct __hip_tex_ret<
+    T,
+    hipReadModeElementType,
+    typename std::enable_if<__hip_is_channel_type<T>::value, bool>::type>
+{
+    using type = T;
+};
+
+template<
+    typename T,
+    unsigned int rank>
+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>
+{
+    using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeElementType>, rank>;
+};
+
+template<typename T>
+struct __hip_tex_ret<
+    T,
+    hipReadModeNormalizedFloat,
+    typename std::enable_if<__hip_is_normalized_channel_type<T>::value, bool>::type>
+{
+    using type = float;
+};
+
+template<
+    typename T,
+    unsigned int rank>
+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>
+{
+    using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeNormalizedFloat>, rank>;
+};
 
 template <typename T, hipTextureReadMode readMode>
 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1Dfetch(texture<T, hipTextureType1D, readMode> t, int x)
@@ -229,20 +315,48 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3DGrad(texture
     return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
 }
 
-template <typename T, hipTextureReadMode readMode>
-struct __hip_tex2dgather_ret {};
+template <
+    typename T,
+    hipTextureReadMode readMode,
+    typename Enable = void>
+struct __hip_tex2dgather_ret
+{
+    static_assert(std::is_same<Enable, void>::value, "Invalid channel type!");
+};
 
-template <typename T, hipTextureReadMode readMode>
-using __hip_tex2dgather_ret_t = typename __hip_tex2dgather_ret<T, readMode>::type;
+template <
+    typename T,
+    hipTextureReadMode readMode>
+using __hip_tex2dgather_ret_t = typename __hip_tex2dgather_ret<T, readMode, bool>::type;
 
 template <typename T>
-struct __hip_tex2dgather_ret<T, hipReadModeElementType> { using type = HIP_vector_type<T, 4>; };
+struct __hip_tex2dgather_ret<
+    T,
+    hipReadModeElementType,
+    typename std::enable_if<__hip_is_channel_type<T>::value, bool>::type>
+{
+    using type = HIP_vector_type<T, 4>;
+};
 
-template<typename T, unsigned int rank>
-struct __hip_tex2dgather_ret<HIP_vector_type<T, rank>, hipReadModeElementType> { using type = HIP_vector_type<T, 4>; };
+template<
+    typename T,
+    unsigned int rank>
+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>
+{
+    using type = HIP_vector_type<T, 4>;
+};
 
 template <typename T>
-struct __hip_tex2dgather_ret<T, hipReadModeNormalizedFloat> { using type = float4; };
+struct __hip_tex2dgather_ret<
+    T,
+    hipReadModeNormalizedFloat,
+    typename std::enable_if<__hip_is_normalized_channel_type<T>::value, bool>::type>
+{
+    using type = float4;
+};
 
 template <typename T, hipTextureReadMode readMode>
 static __forceinline__ __device__ __hip_tex2dgather_ret_t<T, readMode> tex2Dgather(texture<T, hipTextureType2D, readMode> t, float x, float y, int comp=0)