21 inline namespace _V1 {
22 namespace ext::intel::esimd::detail {
27 template <>
struct element_type_traits<tfloat32> {
28 using RawT =
unsigned int;
29 using EnclosingCppT = float;
31 static constexpr
bool use_native_cpp_ops =
false;
32 static constexpr
bool is_floating_point =
true;
37 template <
int N>
struct vector_conversion_traits<tfloat32, N> {
38 using StdT = __cpp_t<tfloat32>;
39 using RawT = __raw_t<tfloat32>;
41 static ESIMD_INLINE vector_type_t<RawT, N>
42 convert_to_raw(vector_type_t<StdT, N> Val) {
43 #ifdef __SYCL_DEVICE_ONLY__
44 vector_type_t<RawT, N> Result = __esimd_tf32_cvt<RawT, StdT, N>(Val);
47 vector_type_t<RawT, N> Output = 0;
49 for (
int i = 0; i < N; i++) {
50 Output[i] = sycl::bit_cast<RawT>(
static_cast<tfloat32
>(Val[i]));
56 static ESIMD_INLINE vector_type_t<StdT, N>
57 convert_to_cpp(vector_type_t<RawT, N> Val) {
58 vector_type_t<StdT, N> Result = sycl::bit_cast<vector_type_t<StdT, N>>(Val);
63 template <>
struct scalar_conversion_traits<tfloat32> {
64 using RawT = __raw_t<tfloat32>;
66 static ESIMD_INLINE RawT bitcast_to_raw(tfloat32 Val) {
67 return sycl::bit_cast<RawT>(Val);
70 static ESIMD_INLINE tfloat32 bitcast_to_wrapper(RawT Val) {
71 return sycl::bit_cast<tfloat32>(Val);
76 inline std::ostream &
operator<<(std::ostream &O, tfloat32
const &rhs) {
77 O << static_cast<float>(rhs);
81 template <>
struct is_esimd_arithmetic_type<tfloat32, void> : std::true_type {};
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)