From 650b6767c42b9aaa6643e06da3043eacf458c982 Mon Sep 17 00:00:00 2001 From: taosang2 Date: Thu, 20 Apr 2023 11:06:50 -0400 Subject: [PATCH] SWDEV-390626 - Fix wrong mapping functions Remove wrong functions __hipMapFromNativeFloat4() __hipMapToNativeFloat4() Replace them with mapFrom() and mapTo() Change-Id: I1692b92d397bfd732d562778f918ebf8ca532e93 --- .../hip/amd_detail/amd_hip_vector_types.h | 143 +++++++++++++++--- .../hip/amd_detail/amd_surface_functions.h | 130 ++-------------- 2 files changed, 138 insertions(+), 135 deletions(-) diff --git a/hipamd/include/hip/amd_detail/amd_hip_vector_types.h b/hipamd/include/hip/amd_detail/amd_hip_vector_types.h index 8215fb02e2..9451ac65eb 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_vector_types.h +++ b/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(u.x); - return t; + return { + static_cast(u.x) + }; } template @@ -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(u.x); - t.y = static_cast(u.y); - return t; + return { + static_cast(u.x), + static_cast(u.y) + }; } template @@ -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(u.x); - t.y = static_cast(u.y); - t.z = static_cast(u.z); - return t; + return { + static_cast(u.x), + static_cast(u.y), + static_cast(u.z) + }; +} + +template +__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(u.x), + static_cast(0), + static_cast(0), + static_cast(0) + }; +} + +template +__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(u.x), + static_cast(u.y), + static_cast(0), + static_cast(0) + }; +} + +template +__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(u.x), + static_cast(u.y), + static_cast(u.z), + static_cast(0) + }; } template @@ -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(u.x); - t.y = static_cast(u.y); - t.z = static_cast(u.z); - t.w = static_cast(u.w); - return t; + return { + static_cast(u.x), + static_cast(u.y), + static_cast(u.z), + static_cast(u.w) + }; } template @@ -2135,13 +2180,73 @@ mapFrom(const U &u) { int4 i4; uint4 u4; } d = { u }; - if(std::is_signed::value) { + if (std::is_signed::value) { return mapElem(d.i4) ; } else { return mapElem(d.u4); } } +template +__HOST_DEVICE__ +__forceinline__ +typename std::enable_if< + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value, const U>::type +mapTo(const T &t) { + union { + U u; + T t; + } d = { 0 }; + d.t = t; + return d.u; +} + +template +__HOST_DEVICE__ +__forceinline__ +typename std::enable_if< + (sizeof(T) == sizeof(typename T::value_type)) || + std::is_same::value || + std::is_same::value || + std::is_same::value, const U>::type +mapTo(const T &t) { + union { + U u; + T t; + } d = { 0 }; + d.t = t; + return d.u; +} + +template +__HOST_DEVICE__ +__forceinline__ +typename std::enable_if< + (sizeof(T) > sizeof(typename T::value_type)) && ( + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value), const U>::type +mapTo(const T &t) { + union { + U u; + int4 i4; + uint4 u4; + } d = { 0 }; + if (std::is_signed::value) { + d.i4 = mapElem(t); + } else { + d.u4 = mapElem(t); + } + return d.u; +} + #else #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ static inline __HOST_DEVICE__ type make_##type(comp x) { \ diff --git a/hipamd/include/hip/amd_detail/amd_surface_functions.h b/hipamd/include/hip/amd_detail/amd_surface_functions.h index bae218ee14..c59995ca39 100644 --- a/hipamd/include/hip/amd_detail/amd_surface_functions.h +++ b/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::value>::type* = nullptr> -static __HOST_DEVICE__ __forceinline__ float4::Native_vec_ __hipMapToNativeFloat4(const T& t) { - float4::Native_vec_ tmp; - tmp.x = static_cast(t); - return tmp; -} - -template < - typename T, - typename std::enable_if::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(t.x); - return tmp; -} - -template < - typename T, - typename std::enable_if::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(t.x); - tmp.y = static_cast(t.y); - return tmp; -} - -template < - typename T, - typename std::enable_if::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(t.x); - tmp.y = static_cast(t.y); - tmp.z = static_cast(t.z); - return tmp; -} - -template < - typename T, - typename std::enable_if::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(t.x); - tmp.y = static_cast(t.y); - tmp.z = static_cast(t.z); - tmp.w = static_cast(t.w); - return tmp; -} - -template -static __HOST_DEVICE__ __forceinline__ -typename std::enable_if::value, const T>::type -__hipMapFromNativeFloat4(const float4::Native_vec_& u) { - T tmp; - tmp = static_cast(u.x); - return tmp; -} - -template -static __HOST_DEVICE__ __forceinline__ -typename std::enable_if::value && sizeof(T) / sizeof(typename T::value_type) == 1, const T>::type -__hipMapFromNativeFloat4(const float4::Native_vec_& u) { - T tmp; - tmp.x = static_cast(u.x); - return tmp; -} - -template -static __HOST_DEVICE__ __forceinline__ -typename std::enable_if::value && sizeof(T) / sizeof(typename T::value_type) == 2, const T>::type -__hipMapFromNativeFloat4(const float4::Native_vec_& u) { - T tmp; - tmp.x = static_cast(u.x); - tmp.y = static_cast(u.y); - return tmp; -} - -template -static __HOST_DEVICE__ __forceinline__ -typename std::enable_if::value && sizeof(T) / sizeof(typename T::value_type) == 3, const T>::type -__hipMapFromNativeFloat4(const float4::Native_vec_& u) { - T tmp; - tmp.x = static_cast(u.x); - tmp.y = static_cast(u.y); - tmp.z = static_cast(u.z); - return tmp; -} - -template -static __HOST_DEVICE__ __forceinline__ -typename std::enable_if::value && sizeof(T) / sizeof(typename T::value_type) == 4, const T>::type -__hipMapFromNativeFloat4(const float4::Native_vec_& u) { - T tmp; - tmp.x = static_cast(u.x); - tmp.y = static_cast(u.y); - tmp.z = static_cast(u.z); - tmp.w = static_cast(u.w); - return tmp; -} - template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::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(tmp); + *data = mapFrom(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(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(tmp); + *data = mapFrom(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(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(tmp); + *data = mapFrom(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(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(tmp); + *data = mapFrom(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(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(tmp); + *data = mapFrom(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(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(tmp); + *data = mapFrom(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(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(tmp); + *data = mapFrom(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(data); __ockl_image_store_lod_CM(i, int2(x, y).data, face, layer, tmp); }