Skip to content

Commit fb66f1b

Browse files
authored
[SYCL] Replace __builtin_bit_cast with sycl::bit_cast in imf headers (#13313)
Signed-off-by: jinge90 <[email protected]>
1 parent b49d27f commit fb66f1b

File tree

1 file changed

+18
-17
lines changed

1 file changed

+18
-17
lines changed

sycl/include/sycl/ext/intel/math.hpp

Lines changed: 18 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -19,6 +19,7 @@ using _iml_half_internal = _Float16;
1919
using _iml_half_internal = uint16_t;
2020
#endif
2121

22+
#include <sycl/bit_cast.hpp>
2223
#include <sycl/builtins.hpp>
2324
#include <sycl/ext/intel/math/imf_fp_conversions.hpp>
2425
#include <sycl/ext/intel/math/imf_half_trivial.hpp>
@@ -82,9 +83,9 @@ std::enable_if_t<std::is_same_v<Tp, double>, double> copysign(Tp x, Tp y) {
8283
template <typename Tp>
8384
std::enable_if_t<std::is_same_v<Tp, sycl::half>, sycl::half> copysign(Tp x,
8485
Tp y) {
85-
_iml_half_internal xi = __builtin_bit_cast(_iml_half_internal, x);
86-
_iml_half_internal yi = __builtin_bit_cast(_iml_half_internal, y);
87-
return __builtin_bit_cast(sycl::half, __imf_copysignf16(xi, yi));
86+
_iml_half_internal xi = sycl::bit_cast<_iml_half_internal>(x);
87+
_iml_half_internal yi = sycl::bit_cast<_iml_half_internal>(y);
88+
return sycl::bit_cast<sycl::half>(__imf_copysignf16(xi, yi));
8889
}
8990

9091
template <typename Tp>
@@ -99,8 +100,8 @@ std::enable_if_t<std::is_same_v<Tp, double>, double> ceil(Tp x) {
99100

100101
template <typename Tp>
101102
std::enable_if_t<std::is_same_v<Tp, sycl::half>, sycl::half> ceil(Tp x) {
102-
_iml_half_internal xi = __builtin_bit_cast(_iml_half_internal, x);
103-
return __builtin_bit_cast(sycl::half, __imf_ceilf16(xi));
103+
_iml_half_internal xi = sycl::bit_cast<_iml_half_internal>(x);
104+
return sycl::bit_cast<sycl::half>(__imf_ceilf16(xi));
104105
}
105106

106107
template <typename Tp>
@@ -120,8 +121,8 @@ std::enable_if_t<std::is_same_v<Tp, double>, double> floor(Tp x) {
120121

121122
template <typename Tp>
122123
std::enable_if_t<std::is_same_v<Tp, sycl::half>, sycl::half> floor(Tp x) {
123-
_iml_half_internal xi = __builtin_bit_cast(_iml_half_internal, x);
124-
return __builtin_bit_cast(sycl::half, __imf_floorf16(xi));
124+
_iml_half_internal xi = sycl::bit_cast<_iml_half_internal>(x);
125+
return sycl::bit_cast<sycl::half>(__imf_floorf16(xi));
125126
}
126127

127128
template <typename Tp>
@@ -141,8 +142,8 @@ std::enable_if_t<std::is_same_v<Tp, double>, double> inv(Tp x) {
141142

142143
template <typename Tp>
143144
std::enable_if_t<std::is_same_v<Tp, sycl::half>, sycl::half> inv(Tp x) {
144-
_iml_half_internal xi = __builtin_bit_cast(_iml_half_internal, x);
145-
return __builtin_bit_cast(sycl::half, __imf_invf16(xi));
145+
_iml_half_internal xi = sycl::bit_cast<_iml_half_internal>(x);
146+
return sycl::bit_cast<sycl::half>(__imf_invf16(xi));
146147
}
147148

148149
template <typename Tp>
@@ -162,8 +163,8 @@ std::enable_if_t<std::is_same_v<Tp, double>, double> rint(Tp x) {
162163

163164
template <typename Tp>
164165
std::enable_if_t<std::is_same_v<Tp, sycl::half>, sycl::half> rint(Tp x) {
165-
_iml_half_internal xi = __builtin_bit_cast(_iml_half_internal, x);
166-
return __builtin_bit_cast(sycl::half, __imf_rintf16(xi));
166+
_iml_half_internal xi = sycl::bit_cast<_iml_half_internal>(x);
167+
return sycl::bit_cast<sycl::half>(__imf_rintf16(xi));
167168
}
168169

169170
template <typename Tp>
@@ -183,8 +184,8 @@ std::enable_if_t<std::is_same_v<Tp, double>, double> sqrt(Tp x) {
183184

184185
template <typename Tp>
185186
std::enable_if_t<std::is_same_v<Tp, sycl::half>, sycl::half> sqrt(Tp x) {
186-
_iml_half_internal xi = __builtin_bit_cast(_iml_half_internal, x);
187-
return __builtin_bit_cast(sycl::half, __imf_sqrtf16(xi));
187+
_iml_half_internal xi = sycl::bit_cast<_iml_half_internal>(x);
188+
return sycl::bit_cast<sycl::half>(__imf_sqrtf16(xi));
188189
}
189190

190191
template <typename Tp>
@@ -204,8 +205,8 @@ std::enable_if_t<std::is_same_v<Tp, double>, double> rsqrt(Tp x) {
204205

205206
template <typename Tp>
206207
std::enable_if_t<std::is_same_v<Tp, sycl::half>, sycl::half> rsqrt(Tp x) {
207-
_iml_half_internal xi = __builtin_bit_cast(_iml_half_internal, x);
208-
return __builtin_bit_cast(sycl::half, __imf_rsqrtf16(xi));
208+
_iml_half_internal xi = sycl::bit_cast<_iml_half_internal>(x);
209+
return sycl::bit_cast<sycl::half>(__imf_rsqrtf16(xi));
209210
}
210211

211212
template <typename Tp>
@@ -225,8 +226,8 @@ std::enable_if_t<std::is_same_v<Tp, double>, double> trunc(Tp x) {
225226

226227
template <typename Tp>
227228
std::enable_if_t<std::is_same_v<Tp, sycl::half>, sycl::half> trunc(Tp x) {
228-
_iml_half_internal xi = __builtin_bit_cast(_iml_half_internal, x);
229-
return __builtin_bit_cast(sycl::half, __imf_truncf16(xi));
229+
_iml_half_internal xi = sycl::bit_cast<_iml_half_internal>(x);
230+
return sycl::bit_cast<sycl::half>(__imf_truncf16(xi));
230231
}
231232

232233
template <typename Tp>

0 commit comments

Comments
 (0)