[SYCL] Rework 'half' implementation in order to remove bunch of workarounds (#1089)

Because of the fact, that `half` type is not a standard C++ type and it is
not supported everywhere, its implementation differs between host and
device: C++ class with overloaded arithmetic operators is used on host
and `_Float16` is used on device side.

Previously, the switch between two version was implemented as
preprocessor macro and having two different types caused some problems
with integration header and unnamed lambda feature, see #185 and
#960.

This patch redesigned `half` implementation in a way, that single
wrapper data type is used as `half` representation on both host and
device sides; differentiation between actual host and device
implementations is done under the hood of this wrapper.

Signed-off-by: Alexey Sachkov <alexey.sachkov@intel.com>

Author: AlexeySachkov

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());
 }
 

clang/test/CodeGenSYCL/half-with-unnamed-lambda.cpp

@@ -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/aliases.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/boolean.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_lists.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/generic_type_traits.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/kernel_desc.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>