21 inline namespace _V1 {
22 namespace ext::intel::esimd::detail {
26 template <>
struct element_type_traits<
bfloat16> {
28 using RawT = uint_type_t<
sizeof(
bfloat16)>;
31 using EnclosingCppT = float;
33 static constexpr
bool use_native_cpp_ops =
false;
34 static constexpr
bool is_floating_point =
true;
37 #ifdef __SYCL_DEVICE_ONLY__
40 using vc_be_bfloat16_raw_t = _Float16;
45 template <
int N>
struct vector_conversion_traits<
bfloat16, N> {
46 using StdT = __cpp_t<bfloat16>;
47 using StdVecT = vector_type_t<StdT, N>;
48 using RawT = __raw_t<bfloat16>;
50 static ESIMD_INLINE vector_type_t<RawT, N>
51 convert_to_raw(vector_type_t<StdT, N> Val) {
52 #ifdef __SYCL_DEVICE_ONLY__
53 using RawVecT = vector_type_t<vc_be_bfloat16_raw_t, N>;
54 RawVecT ConvVal = __esimd_bf_cvt<vc_be_bfloat16_raw_t, StdT, N>(Val);
56 return sycl::bit_cast<vector_type_t<RawT, N>>(ConvVal);
58 __ESIMD_UNSUPPORTED_ON_HOST;
62 static ESIMD_INLINE vector_type_t<StdT, N>
63 convert_to_cpp(vector_type_t<RawT, N> Val) {
64 #ifdef __SYCL_DEVICE_ONLY__
65 using RawVecT = vector_type_t<vc_be_bfloat16_raw_t, N>;
66 RawVecT Bits = sycl::bit_cast<RawVecT>(Val);
67 return __esimd_bf_cvt<StdT, vc_be_bfloat16_raw_t, N>(Bits);
69 __ESIMD_UNSUPPORTED_ON_HOST;
76 template <>
struct scalar_conversion_traits<
bfloat16> {
77 using RawT = __raw_t<bfloat16>;
79 static ESIMD_INLINE RawT bitcast_to_raw(
bfloat16 Val) {
80 return sycl::bit_cast<RawT>(Val);
83 static ESIMD_INLINE
bfloat16 bitcast_to_wrapper(RawT Val) {
84 return sycl::bit_cast<bfloat16>(Val);
93 O << static_cast<float>(rhs);
97 template <>
struct is_esimd_arithmetic_type<
bfloat16, void> : std::true_type {};
sycl::ext::oneapi::bfloat16 bfloat16
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)