SWDEV-390626 - Fix wrong mapping functions

Remove wrong functions
__hipMapFromNativeFloat4()
__hipMapToNativeFloat4()

Replace them with mapFrom() and mapTo()

Change-Id: I1692b92d397bfd732d562778f918ebf8ca532e93

hipamd/include/hip/amd_detail/amd_hip_vector_types.h

@@ -2038,9 +2038,9 @@ 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;
+  return {
+    static_cast<typename T::value_type>(u.x)
+  };
 }
 
 template<typename T, typename U>
@@ -2050,10 +2050,10 @@ 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;
+  return {
+    static_cast<typename T::value_type>(u.x),
+    static_cast<typename T::value_type>(u.y)
+  };
 }
 
 template<typename T, typename U>
@@ -2063,11 +2063,56 @@ 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;
+  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>
@@ -2077,12 +2122,12 @@ 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;
+  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>
@@ -2135,13 +2180,73 @@ mapFrom(const U &u) {
     int4 i4;
     uint4 u4;
   } d = { u };
-  if(std::is_signed<typename T::value_type>::value) {
+  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) {                                         \

hipamd/include/hip/amd_detail/amd_surface_functions.h

@@ -113,108 +113,6 @@ static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format,
     return x = OrderLUT[order] == 3 ? x / OrderLUT[order] : x >> OrderLUT[order];
 }
 
-template <
-    typename T,
-    typename std::enable_if<std::is_scalar<T>::value>::type* = nullptr>
-static __HOST_DEVICE__ __forceinline__ float4::Native_vec_ __hipMapToNativeFloat4(const T& t) {
-    float4::Native_vec_ tmp;
-    tmp.x = static_cast<float>(t);
-    return tmp;
-}
-
-template <
-    typename T,
-    typename std::enable_if<!std::is_scalar<T>::value && sizeof(T) / sizeof(typename T::value_type) == 1>::type* = nullptr>
-static __HOST_DEVICE__ __forceinline__ float4::Native_vec_ __hipMapToNativeFloat4(const T& t) {
-    float4::Native_vec_ tmp;
-    tmp.x = static_cast<float>(t.x);
-    return tmp;
-}
-
-template <
-    typename T,
-    typename std::enable_if<!std::is_scalar<T>::value && sizeof(T) / sizeof(typename T::value_type) == 2>::type* = nullptr>
-static __HOST_DEVICE__ __forceinline__ float4::Native_vec_ __hipMapToNativeFloat4(const T& t) {
-    float4::Native_vec_ tmp;
-    tmp.x = static_cast<float>(t.x);
-    tmp.y = static_cast<float>(t.y);
-    return tmp;
-}
-
-template <
-    typename T,
-    typename std::enable_if<!std::is_scalar<T>::value && sizeof(T) / sizeof(typename T::value_type) == 3>::type* = nullptr>
-static __HOST_DEVICE__ __forceinline__ float4::Native_vec_ __hipMapToNativeFloat4(const T& t) {
-    float4::Native_vec_ tmp;
-    tmp.x = static_cast<float>(t.x);
-    tmp.y = static_cast<float>(t.y);
-    tmp.z = static_cast<float>(t.z);
-    return tmp;
-}
-
-template <
-    typename T,
-    typename std::enable_if<!std::is_scalar<T>::value && sizeof(T) / sizeof(typename T::value_type) == 4>::type* = nullptr>
-static __HOST_DEVICE__ __forceinline__ float4::Native_vec_ __hipMapToNativeFloat4(const T& t) {
-    float4::Native_vec_ tmp;
-    tmp.x = static_cast<float>(t.x);
-    tmp.y = static_cast<float>(t.y);
-    tmp.z = static_cast<float>(t.z);
-    tmp.w = static_cast<float>(t.w);
-    return tmp;
-}
-
-template<typename T>
-static __HOST_DEVICE__ __forceinline__ 
-typename std::enable_if<std::is_scalar<T>::value, const T>::type 
-__hipMapFromNativeFloat4(const float4::Native_vec_& u) {
-    T tmp;
-    tmp = static_cast<T>(u.x);
-    return tmp;
-}
-
-template<typename T>
-static __HOST_DEVICE__ __forceinline__
-typename std::enable_if<!std::is_scalar<T>::value && sizeof(T) / sizeof(typename T::value_type) == 1, const T>::type 
-__hipMapFromNativeFloat4(const float4::Native_vec_& u) {
-    T tmp;
-    tmp.x = static_cast<typename T::value_type>(u.x);
-    return tmp;
-}
-
-template<typename T>
-static __HOST_DEVICE__ __forceinline__
-typename std::enable_if<!std::is_scalar<T>::value && sizeof(T) / sizeof(typename T::value_type) == 2, const T>::type
-__hipMapFromNativeFloat4(const float4::Native_vec_& u) {
-    T tmp;
-    tmp.x = static_cast<typename T::value_type>(u.x);
-    tmp.y = static_cast<typename T::value_type>(u.y);
-    return tmp;
-}
-
-template<typename T>
-static __HOST_DEVICE__ __forceinline__
-typename std::enable_if<!std::is_scalar<T>::value && sizeof(T) / sizeof(typename T::value_type) == 3, const T>::type
-__hipMapFromNativeFloat4(const float4::Native_vec_& u) {
-    T tmp;
-    tmp.x = static_cast<typename T::value_type>(u.x);
-    tmp.y = static_cast<typename T::value_type>(u.y);
-    tmp.z = static_cast<typename T::value_type>(u.z);
-    return tmp;
-}
-
-template<typename T>
-static __HOST_DEVICE__ __forceinline__
-typename std::enable_if<!std::is_scalar<T>::value && sizeof(T) / sizeof(typename T::value_type) == 4, const T>::type
-__hipMapFromNativeFloat4(const float4::Native_vec_& u) {
-    T tmp;
-    tmp.x = static_cast<typename T::value_type>(u.x);
-    tmp.y = static_cast<typename T::value_type>(u.y);
-    tmp.z = static_cast<typename T::value_type>(u.z);
-    tmp.w = static_cast<typename T::value_type>(u.w);
-    return tmp;
-}
-
 template <
     typename T,
     typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
@@ -223,7 +121,7 @@ static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t su
     __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 = __hipMapFromNativeFloat4<T>(tmp);
+    *data = mapFrom<T>(tmp);
 }
 
 template <
@@ -232,7 +130,7 @@ template <
 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 = __hipMapToNativeFloat4(data);
+    auto tmp = mapTo<float4::Native_vec_>(data);
     __ockl_image_store_1D(i, x, tmp);
 }
 
@@ -243,7 +141,7 @@ static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t su
     __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 = __hipMapFromNativeFloat4<T>(tmp);
+    *data = mapFrom<T>(tmp);
 }
 
 template <
@@ -252,7 +150,7 @@ template <
 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 = __hipMapToNativeFloat4(data);
+    auto tmp = mapTo<float4::Native_vec_>(data);
     __ockl_image_store_2D(i, int2(x, y).data, tmp);
 }
 
@@ -263,7 +161,7 @@ static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t su
     __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 = __hipMapFromNativeFloat4<T>(tmp);
+    *data = mapFrom<T>(tmp);
 }
 
 template <
@@ -272,7 +170,7 @@ template <
 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 = __hipMapToNativeFloat4(data);
+    auto tmp = mapTo<float4::Native_vec_>(data);
     __ockl_image_store_3D(i, int4(x, y, z, 0).data, tmp);
 }
 
@@ -283,7 +181,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObje
     __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 = __hipMapFromNativeFloat4<T>(tmp);
+    *data = mapFrom<T>(tmp);
 }
 
 template <
@@ -292,7 +190,7 @@ template <
 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 = __hipMapToNativeFloat4(data);
+    auto tmp = mapTo<float4::Native_vec_>(data);
     __ockl_image_store_lod_1D(i, x, layer, tmp);
 }
 
@@ -303,7 +201,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObje
     __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 = __hipMapFromNativeFloat4<T>(tmp);
+    *data = mapFrom<T>(tmp);
 }
 
 template <
@@ -312,7 +210,7 @@ template <
 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 = __hipMapToNativeFloat4(data);
+    auto tmp = mapTo<float4::Native_vec_>(data);
     __ockl_image_store_lod_2D(i, int2(x, y).data, layer, tmp);
 }
 
@@ -323,7 +221,7 @@ static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject
     __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 = __hipMapFromNativeFloat4<T>(tmp);
+    *data = mapFrom<T>(tmp);
 }
 
 template <
@@ -332,7 +230,7 @@ template <
 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 = __hipMapToNativeFloat4(data);
+    auto tmp = mapTo<float4::Native_vec_>(data);
     __ockl_image_store_CM(i, int2(x, y).data, face, tmp);
 }
 
@@ -344,7 +242,7 @@ static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfac
     __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 = __hipMapFromNativeFloat4<T>(tmp);
+    *data = mapFrom<T>(tmp);
 }
 
 template <
@@ -354,7 +252,7 @@ static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfa
         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 = __hipMapToNativeFloat4(data);
+    auto tmp = mapTo<float4::Native_vec_>(data);
     __ockl_image_store_lod_CM(i, int2(x, y).data, face, layer, tmp);
 }