diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bf16.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bf16.h index a731db3e6e..132abf713d 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bf16.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bf16.h @@ -1660,7 +1660,8 @@ __BF16_DEVICE_STATIC__ __hip_bfloat16 hcos(const __hip_bfloat16 h) { * \brief Calculate exponential of bfloat16 */ __BF16_DEVICE_STATIC__ __hip_bfloat16 hexp(const __hip_bfloat16 h) { - return __float2bfloat16(__ocml_exp_f32(__bfloat162float(h))); + // FIXME: Manual promotion to float unnecessary + return __float2bfloat16(__builtin_elementwise_exp(__bfloat162float(h))); } /** @@ -1676,7 +1677,8 @@ __BF16_DEVICE_STATIC__ __hip_bfloat16 hexp10(const __hip_bfloat16 h) { * \brief Calculate exponential 2 of bfloat16 */ __BF16_DEVICE_STATIC__ __hip_bfloat16 hexp2(const __hip_bfloat16 h) { - return __float2bfloat16(__ocml_exp2_f32(__bfloat162float(h))); + // FIXME: Manual promotion to float unnecessary + return __float2bfloat16(__builtin_elementwise_exp2(__bfloat162float(h))); } /** 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 aaf7b05396..21c69640cb 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 @@ -959,10 +959,10 @@ inline __device__ __half hcos(__half x) { return __half_raw{__ocml_cos_f16(static_cast<__half_raw>(x).data)}; } inline __device__ __half hexp(__half x) { - return __half_raw{__ocml_exp_f16(static_cast<__half_raw>(x).data)}; + return __half_raw{__builtin_elementwise_exp(static_cast<__half_raw>(x).data)}; } inline __device__ __half hexp2(__half x) { - return __half_raw{__ocml_exp2_f16(static_cast<__half_raw>(x).data)}; + return __half_raw{__builtin_elementwise_exp2(static_cast<__half_raw>(x).data)}; } inline __device__ __half hexp10(__half x) { return __half_raw{__ocml_exp10_f16(static_cast<__half_raw>(x).data)}; @@ -1008,8 +1008,12 @@ inline __device__ __half2 h2rint(__half2 x) { } inline __device__ __half2 h2sin(__half2 x) { return __half2{__ocml_sin_2f16(x)}; } inline __device__ __half2 h2cos(__half2 x) { return __half2{__ocml_cos_2f16(x)}; } -inline __device__ __half2 h2exp(__half2 x) { return __half2{__ocml_exp_2f16(x)}; } -inline __device__ __half2 h2exp2(__half2 x) { return __half2{__ocml_exp2_2f16(x)}; } +inline __device__ __half2 h2exp(__half2 x) { + return __half2{__builtin_elementwise_exp(static_cast<__half2_raw>(x).data)}; +} +inline __device__ __half2 h2exp2(__half2 x) { + return __half2{__builtin_elementwise_exp2(static_cast<__half2_raw>(x).data)}; +} inline __device__ __half2 h2exp10(__half2 x) { return __half2{__ocml_exp10_2f16(x)}; } inline __device__ __half2 h2log2(__half2 x) { return __half2{__ocml_log2_2f16(x)}; } inline __device__ __half2 h2log(__half2 x) { return __ocml_log_2f16(x); }