SWDEV-315590 - Fix texture fetch issue

Add mapFrom() to map "float __attribute__((ext_vector_type(4)))"
of output of low-level device function to right vector.

Change-Id: Ifed915c87273b784780b2180320ccb746845dafb


[ROCm/clr commit: 7de480760b]

projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h

@@ -2427,6 +2427,118 @@ 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) {
+  T t;
+  t.x = static_cast<typename T::value_type>(u.x);
+  return t;
+}
+
+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) {
+  T t;
+  t.x = static_cast<typename T::value_type>(u.x);
+  t.y = static_cast<typename T::value_type>(u.y);
+  return t;
+}
+
+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) {
+  T t;
+  t.x = static_cast<typename T::value_type>(u.x);
+  t.y = static_cast<typename T::value_type>(u.y);
+  t.z = static_cast<typename T::value_type>(u.z);
+  return t;
+}
+
+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) {
+  T t;
+  t.x = static_cast<typename T::value_type>(u.x);
+  t.y = static_cast<typename T::value_type>(u.y);
+  t.z = static_cast<typename T::value_type>(u.z);
+  t.w = static_cast<typename T::value_type>(u.w);
+  return t;
+}
+
+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);
+  }
+}
+
 #else
 #define DECLOP_MAKE_ONE_COMPONENT(comp, type)                                                      \
   static inline __HOST_DEVICE__ type make_##type(comp x) {                                         \

projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h

@@ -142,7 +142,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
 {
     TEXTURE_PARAMETERS_INIT;
     auto tmp = __ockl_image_load_1Db(i, x);
-    return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
+    return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
 }
 
 template <typename T, hipTextureReadMode readMode>

projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h

@@ -68,7 +68,7 @@ static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject,
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
     auto tmp = __ockl_image_load_1Db(i, x);
-    return *reinterpret_cast<T*>(&tmp);
+    return mapFrom<T>(tmp);
 }
 
 template <