[SYCL] Rework 'half' implementation in order to remove bunch of workarounds

clang/lib/AST/ItaniumMangle.cpp

@@ -2456,67 +2456,6 @@ static bool isTypeSubstitutable(Qualifiers Quals, const Type *Ty,
   return true;
 }
 
-namespace {
-struct DeclContextDesc {
-  Decl::Kind DeclKind;
-  StringRef Name;
-};
-} // namespace
-
-// For Scopes argument, the only supported Decl::Kind values are:
-// - Namespace
-// - CXXRecord
-// - ClassTemplateSpecialization
-static bool matchQualifiedTypeName(const QualType &Ty,
-                            ArrayRef<DeclContextDesc> Scopes) {
-  // The idea: check the declaration context chain starting from the type
-  // itself. At each step check the context is of expected kind
-  // (namespace) and name.
-  const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
-
-  if (!RecTy)
-    return false; // only classes/structs supported
-  const auto *Ctx = dyn_cast<DeclContext>(RecTy);
-
-  for (const auto &Scope : llvm::reverse(Scopes)) {
-    Decl::Kind DK = Ctx->getDeclKind();
-    StringRef Name = "";
-
-    if (DK != Scope.DeclKind)
-      return false;
-
-    switch (DK) {
-    case Decl::Kind::ClassTemplateSpecialization:
-      // ClassTemplateSpecializationDecl inherits from CXXRecordDecl
-    case Decl::Kind::CXXRecord:
-      Name = cast<CXXRecordDecl>(Ctx)->getName();
-      break;
-    case Decl::Kind::Namespace:
-      Name = cast<NamespaceDecl>(Ctx)->getName();
-      break;
-    default:
-      return false;
-    }
-    if (Name != Scope.Name)
-      return false;
-    Ctx = Ctx->getParent();
-  }
-  return Ctx->isTranslationUnit();
-}
-
-static bool isSYCLHostHalfType(const Type *Ty) {
-  // FIXME: this is not really portable, since the bunch of namespace below
-  // is not specified by the SYCL standard and highly depends on particular
-  // implementation
-  static const std::array<DeclContextDesc, 5> Scopes = {
-      DeclContextDesc{Decl::Kind::Namespace, "cl"},
-      DeclContextDesc{Decl::Kind::Namespace, "sycl"},
-      DeclContextDesc{Decl::Kind::Namespace, "detail"},
-      DeclContextDesc{Decl::Kind::Namespace, "half_impl"},
-      DeclContextDesc{Decl::Kind::CXXRecord, "half"}};
-  return matchQualifiedTypeName(QualType(Ty, 0), Scopes);
-}
-
 void CXXNameMangler::mangleType(QualType T) {
   // If our type is instantiation-dependent but not dependent, we mangle
   // it as it was written in the source, removing any top-level sugar.
@@ -2576,11 +2515,6 @@ void CXXNameMangler::mangleType(QualType T) {
 
   bool isSubstitutable =
     isTypeSubstitutable(quals, ty, Context.getASTContext());
-  if (Context.isUniqueNameMangler() && isSYCLHostHalfType(ty)) {
-    // Set isSubstitutable to false for cl::sycl::detail::half_impl::half
-    // to achieve the same mangling for other components
-    isSubstitutable = false;
-  }
   if (isSubstitutable && mangleSubstitution(T))
     return;
 
@@ -3057,11 +2991,6 @@ void CXXNameMangler::mangleType(const RecordType *T) {
   mangleType(static_cast<const TagType*>(T));
 }
 void CXXNameMangler::mangleType(const TagType *T) {
-  if (Context.isUniqueNameMangler() && isSYCLHostHalfType(T)) {
-    // Mangle cl::sycl::detail::half_imple::half as _Float16
-    mangleType(Context.getASTContext().Float16Ty);
-    return;
-  }
   mangleName(T->getDecl());
 }
 

sycl/include/CL/sycl/aliases.hpp

@@ -24,11 +24,11 @@ class half;
 } // namespace sycl
 } // namespace cl
 
-#ifdef __SYCL_DEVICE_ONLY__
-using half = _Float16;
-#else
+// FIXME: line below exports 'half' into global namespace, which seems incorrect
+// However, SYCL 1.2.1 spec considers 'half' to be a fundamental C++ data type
+// which doesn't exist within the 'cl::sycl' namespace.
+// Related spec issue: KhronosGroup/SYCL-Docs#40
 using half = cl::sycl::detail::half_impl::half;
-#endif
 
 #define MAKE_VECTOR_ALIAS(ALIAS, TYPE, N)                                      \
   using ALIAS##N = cl::sycl::vec<TYPE, N>;
@@ -80,7 +80,8 @@ using ulong = unsigned long;
 using longlong = long long;
 using ulonglong = unsigned long long;
 // TODO cl::sycl::half is not in SYCL specification, but is used by Khronos CTS.
-using half = half;
+// Related tests issue: KhronosGroup/SYCL-CTS#37
+using half = cl::sycl::detail::half_impl::half;
 using cl_bool = bool;
 using cl_char = std::int8_t;
 using cl_uchar = std::uint8_t;

sycl/include/CL/sycl/detail/boolean.hpp

@@ -99,7 +99,7 @@ template <int N> struct Boolean {
 
 private:
   template <int Num> friend struct Assigner;
-  alignas(VectorAlignment<bool, N>::value) DataType value;
+  alignas(detail::vector_alignment<bool, N>::value) DataType value;
 };
 
 template <> struct Boolean<1> {

sycl/include/CL/sycl/detail/generic_type_lists.hpp

@@ -11,15 +11,20 @@
 #include <CL/sycl/access/access.hpp>
 #include <CL/sycl/detail/stl_type_traits.hpp>
 #include <CL/sycl/detail/type_list.hpp>
-#include <CL/sycl/half_type.hpp>
 
 // Generic type name description, which serves as a description for all valid
 // types of parameters to kernel functions
 
-// Forward declaration
+// Forward declarations
 __SYCL_INLINE namespace cl {
 namespace sycl {
 template <typename T, int N> class vec;
+namespace detail {
+namespace half_impl {
+class half;
+}
+} // namespace detail
+using half = detail::half_impl::half;
 } // namespace sycl
 } // namespace cl
 

sycl/include/CL/sycl/detail/generic_type_traits.hpp

@@ -232,6 +232,32 @@ using is_genptr = bool_constant<
 
 template <typename T> using is_nan_type = is_contained<T, gtl::nan_list>;
 
+// nan_types
+template <typename T, typename Enable = void> struct nan_types;
+
+template <typename T>
+struct nan_types<
+    T, enable_if_t<is_contained<T, gtl::unsigned_short_list>::value, T>> {
+  using ret_type = change_base_type_t<T, half>;
+  using arg_type = find_same_size_type_t<gtl::scalar_unsigned_short_list, half>;
+};
+
+template <typename T>
+struct nan_types<
+    T, enable_if_t<is_contained<T, gtl::unsigned_int_list>::value, T>> {
+  using ret_type = change_base_type_t<T, float>;
+  using arg_type = find_same_size_type_t<gtl::scalar_unsigned_int_list, float>;
+};
+
+template <typename T>
+struct nan_types<
+    T,
+    enable_if_t<is_contained<T, gtl::unsigned_long_integer_list>::value, T>> {
+  using ret_type = change_base_type_t<T, double>;
+  using arg_type =
+      find_same_size_type_t<gtl::scalar_unsigned_long_integer_list, double>;
+};
+
 template <typename T> using nan_return_t = typename nan_types<T, T>::ret_type;
 
 template <typename T>
@@ -364,10 +390,14 @@ using select_cl_scalar_intergal_t =
 // select_cl_scalar_t picks corresponding cl_* type for input
 // scalar T or returns T if T is not scalar.
 template <typename T>
-using select_cl_scalar_t =
-    conditional_t<std::is_integral<T>::value, select_cl_scalar_intergal_t<T>,
-                  conditional_t<std::is_floating_point<T>::value,
-                                select_cl_scalar_float_t<T>, T>>;
+using select_cl_scalar_t = conditional_t<
+    std::is_integral<T>::value, select_cl_scalar_intergal_t<T>,
+    conditional_t<
+        std::is_floating_point<T>::value, select_cl_scalar_float_t<T>,
+        // half is a special case: it is implemented differently on host and
+        // device and therefore, might lower to different types
+        conditional_t<std::is_same<T, half>::value,
+                      cl::sycl::detail::half_impl::BIsRepresentationT, T>>>;
 
 // select_cl_vector_or_scalar does cl_* type selection for element type of
 // a vector type T and does scalar type substitution.  If T is not
@@ -378,7 +408,13 @@ template <typename T>
 struct select_cl_vector_or_scalar<
     T, typename std::enable_if<is_vgentype<T>::value>::type> {
   using type =
-      vec<select_cl_scalar_t<typename T::element_type>, T::get_count()>;
+      // select_cl_scalar_t returns _Float16, so, we try to instantiate vec
+      // class with _Float16 DataType, which is not expected there
+      // So, leave vector<half, N> as-is
+      vec<conditional_t<std::is_same<typename T::element_type, half>::value,
+                        typename T::element_type,
+                        select_cl_scalar_t<typename T::element_type>>,
+          T::get_count()>;
 };
 
 template <typename T>

sycl/include/CL/sycl/detail/kernel_desc.hpp

@@ -15,22 +15,10 @@ __SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
-namespace half_impl {
-
-class half;
-// Half type is defined as _Float16 on device and as manually implemented half
-// type on host. Integration header is generated by device compiler so it sees
-// half type as _Float16 and it will add _Float16 to integration header if it
-// is used in kernel name template parameters. To avoid errors in host
-// compilation we remove _Float16 from integration header using following macro.
-// Same thing goes about bool type which is defined as _Bool.
 #ifndef __SYCL_DEVICE_ONLY__
-#define _Float16 cl::sycl::detail::half_impl::half
 #define _Bool bool
 #endif
 
-} // namespace half_impl
-
 // kernel parameter kinds
 enum class kernel_param_kind_t {
   kind_accessor,

sycl/include/CL/sycl/detail/type_traits.hpp

@@ -12,12 +12,17 @@
 #include <CL/sycl/detail/generic_type_lists.hpp>
 #include <CL/sycl/detail/stl_type_traits.hpp>
 #include <CL/sycl/detail/type_list.hpp>
-#include <CL/sycl/half_type.hpp>
 
 #include <type_traits>
 
 __SYCL_INLINE namespace cl {
 namespace sycl {
+namespace detail {
+namespace half_impl {
+class half;
+}
+} // namespace detail
+using half = detail::half_impl::half;
 
 // Forward declaration
 template <typename ElementType, access::address_space Space> class multi_ptr;
@@ -38,6 +43,16 @@ struct vector_size_impl<vec<T, N>> : int_constant<N> {};
 template <typename T>
 struct vector_size : vector_size_impl<remove_cv_t<remove_reference_t<T>>> {};
 
+// 4.10.2.6 Memory layout and alignment
+template <typename T, int N>
+struct vector_alignment_impl
+    : conditional_t<N == 3, int_constant<sizeof(T) * 4>,
+                    int_constant<sizeof(T) * N>> {};
+
+template <typename T, int N>
+struct vector_alignment
+    : vector_alignment_impl<remove_cv_t<remove_reference_t<T>>, N> {};
+
 // vector_element
 template <typename T> struct vector_element_impl;
 template <typename T>
@@ -233,32 +248,6 @@ template <typename T, int N, typename TL> struct make_type_impl<vec<T, N>, TL> {
 template <typename T, typename TL>
 using make_type_t = typename make_type_impl<T, TL>::type;
 
-// nan_types
-template <typename T, typename Enable = void> struct nan_types;
-
-template <typename T>
-struct nan_types<
-    T, enable_if_t<is_contained<T, gtl::unsigned_short_list>::value, T>> {
-  using ret_type = change_base_type_t<T, half>;
-  using arg_type = find_same_size_type_t<gtl::scalar_unsigned_short_list, half>;
-};
-
-template <typename T>
-struct nan_types<
-    T, enable_if_t<is_contained<T, gtl::unsigned_int_list>::value, T>> {
-  using ret_type = change_base_type_t<T, float>;
-  using arg_type = find_same_size_type_t<gtl::scalar_unsigned_int_list, float>;
-};
-
-template <typename T>
-struct nan_types<
-    T,
-    enable_if_t<is_contained<T, gtl::unsigned_long_integer_list>::value, T>> {
-  using ret_type = change_base_type_t<T, double>;
-  using arg_type =
-      find_same_size_type_t<gtl::scalar_unsigned_long_integer_list, double>;
-};
-
 // make_larger_t
 template <typename T, typename Enable = void> struct make_larger_impl;
 template <typename T>