diff --git a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h index c3a790751b..24728550ad 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h @@ -125,22 +125,22 @@ uint64_t __make_mantissa(const char* tagp) __DEVICE__ inline int amd_mixed_dot(short2 a, short2 b, int c, bool saturate) { - return __ockl_sdot2(a, b, c, saturate); + return __ockl_sdot2(a.data, b.data, c, saturate); } __DEVICE__ inline uint amd_mixed_dot(ushort2 a, ushort2 b, uint c, bool saturate) { - return __ockl_udot2(a, b, c, saturate); + return __ockl_udot2(a.data, b.data, c, saturate); } __DEVICE__ inline int amd_mixed_dot(char4 a, char4 b, int c, bool saturate) { - return __ockl_sdot4(a, b, c, saturate); + return __ockl_sdot4(a.data, b.data, c, saturate); } __DEVICE__ inline uint amd_mixed_dot(uchar4 a, uchar4 b, uint c, bool saturate) { - return __ockl_udot4(a, b, c, saturate); + return __ockl_udot4(a.data, b.data, c, saturate); } __DEVICE__ inline diff --git a/projects/clr/hipamd/include/hip/hcc_detail/math_fwd.h b/projects/clr/hipamd/include/hip/hcc_detail/math_fwd.h index cdf8e78da4..3021f92d13 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/math_fwd.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/math_fwd.h @@ -32,22 +32,39 @@ THE SOFTWARE. #if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__) __device__ __attribute__((const)) -int __ockl_sdot2(short2 a, short2 b, int c, bool s); +int __ockl_sdot2( + HIP_vector_base::Native_vec_, + HIP_vector_base::Native_vec_, + int, bool); + __device__ __attribute__((const)) -unsigned int __ockl_udot2(ushort2 a, ushort2 b, unsigned int c, bool s); +unsigned int __ockl_udot2( + HIP_vector_base::Native_vec_, + HIP_vector_base::Native_vec_, + unsigned int, bool); + __device__ __attribute__((const)) -int __ockl_sdot4(char4 a, char4 b, int c, bool s); +int __ockl_sdot4( + HIP_vector_base::Native_vec_, + HIP_vector_base::Native_vec_, + int, bool); + __device__ __attribute__((const)) -unsigned int __ockl_udot4(uchar4 a, uchar4 b, unsigned int c, bool s); +unsigned int __ockl_udot4( + HIP_vector_base::Native_vec_, + HIP_vector_base::Native_vec_, + unsigned int, bool); + __device__ __attribute__((const)) -int __ockl_sdot8(int a, int b, int c, bool s); +int __ockl_sdot8(int, int, int, bool); + __device__ __attribute__((const)) -unsigned int __ockl_udot8(unsigned int a, unsigned int b, unsigned int c, bool s); +unsigned int __ockl_udot8(unsigned int, unsigned int, unsigned int, bool); #endif // BEGIN FLOAT