From 9c28f156e43537aadefad743395230f7dc633840 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Thu, 13 Jul 2023 23:03:11 +0100 Subject: [PATCH] SWDEV-411156 - Use _Float16 instead of unsigned short values Change-Id: I23387bf27b2bd70fb47a3a36ffaf114e6c03e0bc [ROCm/clr commit: 295963fe7af03f51013ccacc8c290eda481a71c3] --- projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h index 268238531d..4c6d1cd181 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h @@ -1617,7 +1617,7 @@ THE SOFTWARE. __half hrcp(__half x) { return __half_raw{ - static_cast<_Float16>(__builtin_amdgcn_rcph(static_cast<__half_raw>(x).data))}; + static_cast<_Float16>(1.0f) /static_cast<__half_raw>(x).data}; } inline __device__ @@ -1721,8 +1721,8 @@ THE SOFTWARE. inline __HOST_DEVICE__ __half2 h2rcp(__half2 x) { - return _Float16_2{static_cast<_Float16>(__builtin_amdgcn_rcph(x.x)), - static_cast<_Float16>(__builtin_amdgcn_rcph(x.y))}; + return _Float16_2{ + _Float16_2{static_cast<_Float16>(1.0f), static_cast<_Float16>(1.0f)} / x.data}; } inline __HOST_DEVICE__