From 9f7641070bd8ed84b3cd230d6368048d5f45b7fe Mon Sep 17 00:00:00 2001
From: jinge90 <ge.jin@intel.com>
Date: Sun, 29 Jan 2023 10:55:15 +0800
Subject: [PATCH 1/2] [SYCL] Add _imf_inv functions to sycl intel math

---
 sycl/include/sycl/ext/intel/math.hpp | 20 ++++++++++++++++++++
 1 file changed, 20 insertions(+)

diff --git a/sycl/include/sycl/ext/intel/math.hpp b/sycl/include/sycl/ext/intel/math.hpp
index d0c74a8319c7f..f5af93b1525e8 100644
--- a/sycl/include/sycl/ext/intel/math.hpp
+++ b/sycl/include/sycl/ext/intel/math.hpp
@@ -36,6 +36,9 @@ double __imf_floor(double);
 _iml_half_internal __imf_floorf16(_iml_half_internal);
 float __imf_rintf(float);
 double __imf_rint(double);
+_iml_half_internal __imf_invf16(_iml_half_internal);
+float __imf_invf(float);
+double __imf_inv(double);
 _iml_half_internal __imf_rintf16(_iml_half_internal);
 float __imf_sqrtf(float);
 double __imf_sqrt(double);
@@ -118,6 +121,23 @@ sycl::half2 floor(sycl::half2 x) {
   return sycl::half2{floor(x.s0()), floor(x.s1())};
 }
 
+template <typename Tp>
+std::enable_if_t<std::is_same_v<Tp, float>, float> inv(Tp x) {
+  return __imf_invf(x);
+}
+
+template <typename Tp>
+std::enable_if_t<std::is_same_v<Tp, double>, double> inv(Tp x) {
+  return __imf_inv(x);
+}
+
+template <typename Tp>
+std::enable_if_t<std::is_same_v<Tp, sycl::half>, sycl::half> inv(Tp x) {
+  _iml_half_internal xi = __builtin_bit_cast(_iml_half_internal, x);
+  return __builtin_bit_cast(sycl::half, __imf_invf16(xi));
+}
+
+sycl::half2 inv(sycl::half2 x) { return sycl::half2{inv(x.s0()), inv(x.s1())}; }
 template <typename Tp>
 std::enable_if_t<std::is_same_v<Tp, float>, float> rint(Tp x) {
   return __imf_rintf(x);

From 5cb37f31d293564c76cc360a43535ae2fbcf853b Mon Sep 17 00:00:00 2001
From: jinge90 <ge.jin@intel.com>
Date: Sun, 29 Jan 2023 11:29:10 +0800
Subject: [PATCH 2/2] fix format

---
 sycl/include/sycl/ext/intel/math.hpp | 1 +
 1 file changed, 1 insertion(+)

diff --git a/sycl/include/sycl/ext/intel/math.hpp b/sycl/include/sycl/ext/intel/math.hpp
index f5af93b1525e8..8a6913945083b 100644
--- a/sycl/include/sycl/ext/intel/math.hpp
+++ b/sycl/include/sycl/ext/intel/math.hpp
@@ -138,6 +138,7 @@ std::enable_if_t<std::is_same_v<Tp, sycl::half>, sycl::half> inv(Tp x) {
 }
 
 sycl::half2 inv(sycl::half2 x) { return sycl::half2{inv(x.s0()), inv(x.s1())}; }
+
 template <typename Tp>
 std::enable_if_t<std::is_same_v<Tp, float>, float> rint(Tp x) {
   return __imf_rintf(x);