20 inline namespace _V1 {
21 namespace ext::intel::esimd::detail {
24 #ifdef __SYCL_DEVICE_ONLY__
31 using half_enclosing_cpp_type = half_raw_type;
33 using half_raw_type = uint16_t;
34 using half_enclosing_cpp_type = float;
37 template <>
struct element_type_traits<
sycl::
half> {
38 using RawT = half_raw_type;
39 using EnclosingCppT = half_enclosing_cpp_type;
40 #ifdef __SYCL_DEVICE_ONLY__
43 static constexpr
bool use_native_cpp_ops =
true;
47 static constexpr
bool use_native_cpp_ops =
false;
50 static constexpr
bool is_floating_point =
true;
55 template <
int N>
struct vector_conversion_traits<
sycl::
half, N> {
56 using StdT = half_enclosing_cpp_type;
57 using RawT = half_raw_type;
59 static ESIMD_INLINE vector_type_t<RawT, N>
60 convert_to_raw(vector_type_t<StdT, N> Val)
61 #ifdef __SYCL_DEVICE_ONLY__
66 __ESIMD_UNSUPPORTED_ON_HOST;
70 static ESIMD_INLINE vector_type_t<StdT, N>
71 convert_to_cpp(vector_type_t<RawT, N> Val)
72 #ifdef __SYCL_DEVICE_ONLY__
77 __ESIMD_UNSUPPORTED_ON_HOST;
86 class WrapperElementTypeProxy {
88 static ESIMD_INLINE half_raw_type bitcast_to_raw_scalar(
sycl::half Val) {
89 #ifdef __SYCL_DEVICE_ONLY__
96 static ESIMD_INLINE
sycl::half bitcast_to_wrapper_scalar(half_raw_type Val) {
97 #ifndef __SYCL_DEVICE_ONLY__
98 __ESIMD_UNSUPPORTED_ON_HOST;
107 template <>
struct scalar_conversion_traits<
sycl::
half> {
108 using RawT = half_raw_type;
110 static ESIMD_INLINE RawT bitcast_to_raw(
sycl::half Val) {
111 return WrapperElementTypeProxy::bitcast_to_raw_scalar(Val);
114 static ESIMD_INLINE
sycl::half bitcast_to_wrapper(RawT Val) {
115 return WrapperElementTypeProxy::bitcast_to_wrapper_scalar(Val);
119 #ifdef __SYCL_DEVICE_ONLY__
121 struct is_esimd_arithmetic_type<half_raw_type, void> : std::true_type {};
125 struct is_esimd_arithmetic_type<
sycl::
half, void> : std::true_type {};
129 O << static_cast<float>(rhs);
134 float ValFloat = 0.0f;
detail::host_half_impl::half StorageT
auto operator>>(const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &RHS)
auto operator<<(const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &RHS)
sycl::detail::half_impl::half half