Skip to content

Commit 4d2198e

Browse files
[NFCI][SYCL] Don't use builtins' internals in ext_oneapi_native_math impl (#13154)
Old builtins implementation is soon to be dropped during ABI breaking window. Ensure that this extension implementation doesn't rely on it.
1 parent f381d5c commit 4d2198e

File tree

1 file changed

+6
-8
lines changed

1 file changed

+6
-8
lines changed

sycl/include/sycl/ext/oneapi/experimental/builtins.hpp

Lines changed: 6 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -98,7 +98,7 @@ tanh(T x) __NOEXC {
9898
return sycl::detail::convertFromOpenCLTypeFor<T>(
9999
__clc_native_tanh(sycl::detail::convertToOpenCLType(x)));
100100
#else
101-
return __sycl_std::__invoke_tanh<T>(x);
101+
return sycl::tanh(x);
102102
#endif
103103
}
104104

@@ -120,16 +120,15 @@ inline __SYCL_ALWAYS_INLINE
120120
#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
121121
auto partial_res = native::tanh(sycl::detail::to_vec2(x, i * 2));
122122
#else
123-
auto partial_res = __sycl_std::__invoke_tanh<sycl::vec<T, 2>>(
124-
sycl::detail::to_vec2(x, i * 2));
123+
auto partial_res = sycl::tanh(sycl::detail::to_vec2(x, i * 2));
125124
#endif
126125
sycl::detail::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>));
127126
}
128127
if (N % 2) {
129128
#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
130129
res[N - 1] = native::tanh(x[N - 1]);
131130
#else
132-
res[N - 1] = __sycl_std::__invoke_tanh<T>(x[N - 1]);
131+
res[N - 1] = sycl::tanh(x[N - 1]);
133132
#endif
134133
}
135134

@@ -147,7 +146,7 @@ inline __SYCL_ALWAYS_INLINE
147146
return sycl::detail::convertFromOpenCLTypeFor<T>(
148147
__clc_native_exp2(sycl::detail::convertToOpenCLType(x)));
149148
#else
150-
return __sycl_std::__invoke_exp2<T>(x);
149+
return sycl::exp2(x);
151150
#endif
152151
}
153152

@@ -162,16 +161,15 @@ exp2(sycl::marray<half, N> x) __NOEXC {
162161
#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
163162
auto partial_res = native::exp2(sycl::detail::to_vec2(x, i * 2));
164163
#else
165-
auto partial_res = __sycl_std::__invoke_exp2<sycl::vec<half, 2>>(
166-
sycl::detail::to_vec2(x, i * 2));
164+
auto partial_res = sycl::exp2(sycl::detail::to_vec2(x, i * 2));
167165
#endif
168166
sycl::detail::memcpy(&res[i * 2], &partial_res, sizeof(vec<half, 2>));
169167
}
170168
if (N % 2) {
171169
#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
172170
res[N - 1] = native::exp2(x[N - 1]);
173171
#else
174-
res[N - 1] = __sycl_std::__invoke_exp2<half>(x[N - 1]);
172+
res[N - 1] = sycl::exp2(x[N - 1]);
175173
#endif
176174
}
177175
return res;

0 commit comments

Comments
 (0)