SWDEV-548892 - Stop using ocml exp and exp2 functions (#2032)

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)));
 }
 
 /**

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); }