[ESIMD] Fix specialization/instantiation order in traits infra. (#6677)

This allows to avoid errors on Linux, where compiler picks up
instatiation of wrong specialization of elem_type_traits (the default
unusable one) when instatiating __raw_t alias.

Also do not use specialization of elem_type_traits (e.g. __raw_t) when
definition of the trait is not complete.

Make WrapperElementTypeProxy specific to sycl::half trait, because it
won't be needed for other types traits.

Author: kbobrovs

sycl/include/sycl/ext/intel/esimd/detail/elem_type_traits.hpp

@@ -153,14 +153,12 @@ struct element_type_traits<T, std::enable_if_t<is_vectorizable_v<T>>> {
 
 // ------------------- Useful meta-functions and declarations
 
+template <class T> using __raw_t = typename element_type_traits<T>::RawT;
 template <class T>
-using __raw_t = typename __ESIMD_DNS::element_type_traits<T>::RawT;
-template <class T>
-using __cpp_t = typename __ESIMD_DNS::element_type_traits<T>::EnclosingCppT;
+using __cpp_t = typename element_type_traits<T>::EnclosingCppT;
 
 template <class T, int N>
-using __raw_vec_t =
-    vector_type_t<typename __ESIMD_DNS::element_type_traits<T>::RawT, N>;
+using __raw_vec_t = vector_type_t<typename element_type_traits<T>::RawT, N>;
 
 // Note: using RawVecT in comparison result type calculation does *not* mean
 // the comparison is actually performed on the raw types.
@@ -602,22 +600,6 @@ vector_comparison_op_traits<Op, WrapperT, N>::impl(__raw_vec_t<WrapperT, N> X,
       vector_comparison_op_default<Op, T1, N>(X1, Y1));
 }
 
-// Proxy class to access bit representation of a wrapper type both on host and
-// device. Declared as friend to the wrapper types (e.g. sycl::half).
-// Specific type traits implementations (scalar_conversion_traits) can use
-// concrete wrapper type specializations of the static functions in this class
-// to access private fields in the wrapper type (e.g. sycl::half).
-// TODO add this functionality to sycl type implementation? With C++20,
-// std::bit_cast should be a good replacement.
-class WrapperElementTypeProxy {
-public:
-  template <class WrapperT>
-  static inline __raw_t<WrapperT> bitcast_to_raw_scalar(WrapperT Val);
-
-  template <class WrapperT>
-  static inline WrapperT bitcast_to_wrapper_scalar(__raw_t<WrapperT> Val);
-};
-
 // "Generic" version of std::is_floating_point_v which returns "true" also for
 // the wrapper floating-point types such as sycl::half.
 template <typename T>

sycl/include/sycl/ext/intel/esimd/detail/half_type_traits.hpp

@@ -20,22 +20,28 @@ namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
 namespace ext::intel::esimd::detail {
 
-template <class T>
-struct element_type_traits<T, std::enable_if_t<std::is_same_v<T, sycl::half>>> {
-  // Can't use sycl::detail::half_impl::StorageT as RawT for both host and
-  // device as it still maps to struct on/ host (even though the struct is a
-  // trivial wrapper around uint16_t), and for ESIMD we need a type which can be
-  // an element of clang vector.
+// Standalone definitions to use w/o instantiating element_type_traits.
+#ifdef __SYCL_DEVICE_ONLY__
+// Can't use sycl::detail::half_impl::StorageT as RawT for both host and
+// device as it still maps to struct on/ host (even though the struct is a
+// trivial wrapper around uint16_t), and for ESIMD we need a type which can be
+// an element of clang vector.
+using half_raw_type = sycl::detail::half_impl::StorageT;
+// On device, _Float16 is native Cpp type, so it is the enclosing C++ type
+using half_enclosing_cpp_type = half_raw_type;
+#else
+using half_raw_type = uint16_t;
+using half_enclosing_cpp_type = float;
+#endif // __SYCL_DEVICE_ONLY__
+
+template <> struct element_type_traits<sycl::half> {
+  using RawT = half_raw_type;
+  using EnclosingCppT = half_enclosing_cpp_type;
 #ifdef __SYCL_DEVICE_ONLY__
-  using RawT = sycl::detail::half_impl::StorageT;
-  // On device, _Float16 is native Cpp type, so it is the enclosing C++ type
-  using EnclosingCppT = RawT;
   // On device, operations on half are translated to operations on _Float16,
   // which is natively supported by the device compiler
   static inline constexpr bool use_native_cpp_ops = true;
 #else
-  using RawT = uint16_t;
-  using EnclosingCppT = float;
   // On host, we can't use native Cpp '+', '-' etc. over uint16_t to emulate the
   // operations on half type.
   static inline constexpr bool use_native_cpp_ops = false;
@@ -47,8 +53,8 @@ struct element_type_traits<T, std::enable_if_t<std::is_same_v<T, sycl::half>>> {
 // ------------------- Type conversion traits
 
 template <int N> struct vector_conversion_traits<sycl::half, N> {
-  using StdT = __cpp_t<sycl::half>;
-  using RawT = __raw_t<sycl::half>;
+  using StdT = half_enclosing_cpp_type;
+  using RawT = half_raw_type;
 
   static ESIMD_INLINE vector_type_t<RawT, N>
   convert_to_raw(vector_type_t<StdT, N> Val)
@@ -57,7 +63,7 @@ template <int N> struct vector_conversion_traits<sycl::half, N> {
       ;
 #else
   {
-    vector_type_t<__raw_t<sycl::half>, N> Output = 0;
+    vector_type_t<half_raw_type, N> Output = 0;
 
     for (int i = 0; i < N; i += 1) {
       // 1. Convert Val[i] to float (x) using c++ static_cast
@@ -89,46 +95,49 @@ template <int N> struct vector_conversion_traits<sycl::half, N> {
 #endif // __SYCL_DEVICE_ONLY__
 };
 
-// WrapperElementTypeProxy (a friend of sycl::half) must be used to access
-// private fields of the sycl::half.
-template <>
-ESIMD_INLINE __raw_t<sycl::half>
-WrapperElementTypeProxy::bitcast_to_raw_scalar<sycl::half>(sycl::half Val) {
+// Proxy class to access bit representation of a wrapper type both on host and
+// device. Declared as friend to the wrapper types (e.g. sycl::half).
+// Specific type traits implementations (scalar_conversion_traits) can use
+// concrete wrapper type specializations of the static functions in this class
+// to access private fields in the wrapper type (e.g. sycl::half).
+// TODO add this functionality to sycl type implementation? With C++20,
+// std::bit_cast should be a good replacement.
+class WrapperElementTypeProxy {
+public:
+  static ESIMD_INLINE half_raw_type bitcast_to_raw_scalar(sycl::half Val) {
 #ifdef __SYCL_DEVICE_ONLY__
-  return Val.Data;
+    return Val.Data;
 #else
-  return Val.Data.Buf;
+    return Val.Data.Buf;
 #endif // __SYCL_DEVICE_ONLY__
-}
+  }
 
-template <>
-ESIMD_INLINE sycl::half
-WrapperElementTypeProxy::bitcast_to_wrapper_scalar<sycl::half>(
-    __raw_t<sycl::half> Val) {
+  static ESIMD_INLINE sycl::half bitcast_to_wrapper_scalar(half_raw_type Val) {
 #ifndef __SYCL_DEVICE_ONLY__
-  return sycl::half(::sycl::detail::host_half_impl::half(Val));
+    return sycl::half(::sycl::detail::host_half_impl::half(Val));
 #else
-  sycl::half Res;
-  Res.Data = Val;
-  return Res;
+    sycl::half Res;
+    Res.Data = Val;
+    return Res;
 #endif // __SYCL_DEVICE_ONLY__
-}
+  }
+};
 
 template <> struct scalar_conversion_traits<sycl::half> {
-  using RawT = __raw_t<sycl::half>;
+  using RawT = half_raw_type;
 
   static ESIMD_INLINE RawT bitcast_to_raw(sycl::half Val) {
-    return WrapperElementTypeProxy::bitcast_to_raw_scalar<sycl::half>(Val);
+    return WrapperElementTypeProxy::bitcast_to_raw_scalar(Val);
   }
 
   static ESIMD_INLINE sycl::half bitcast_to_wrapper(RawT Val) {
-    return WrapperElementTypeProxy::bitcast_to_wrapper_scalar<sycl::half>(Val);
+    return WrapperElementTypeProxy::bitcast_to_wrapper_scalar(Val);
   }
 };
 
 #ifdef __SYCL_DEVICE_ONLY__
 template <>
-struct is_esimd_arithmetic_type<__raw_t<sycl::half>, void> : std::true_type {};
+struct is_esimd_arithmetic_type<half_raw_type, void> : std::true_type {};
 #endif // __SYCL_DEVICE_ONLY__
 
 // Misc