Skip to content

[SYCL] Add sycl_ext_intel_math_inv #8138

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 4 commits into from
Feb 9, 2023
Merged

[SYCL] Add sycl_ext_intel_math_inv #8138

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
9f76410
[SYCL] Add _imf_inv functions to sycl intel math
jinge90 Jan 29, 2023
5cb37f3
fix format
jinge90 Jan 29, 2023
d5bbee7
Merge remote-tracking branch 'origin/sycl' into add_rcp_to_imf
jinge90 Feb 3, 2023
0860f2c
Merge remote-tracking branch 'origin/sycl' into add_rcp_to_imf
jinge90 Feb 7, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
21 changes: 21 additions & 0 deletions sycl/include/sycl/ext/intel/math.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -118,6 +121,24 @@ 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);
Expand Down