Merge from 'sycl' to 'sycl-web' (#2)

Author: iclsrc

clang/include/clang/Driver/Options.td

@@ -1818,7 +1818,7 @@ def fstrict_overflow : Flag<["-"], "fstrict-overflow">, Group<f_Group>;
 def fintelfpga : Flag<["-"], "fintelfpga">, Group<f_Group>,
   Flags<[CC1Option, CoreOption]>, HelpText<"Perform ahead of time compilation for FPGA">;
 def fsycl : Flag<["-"], "fsycl">, Group<f_Group>, Flags<[CC1Option, NoArgumentUnused, CoreOption]>,
-  HelpText<"generate SYCL code.">;
+  HelpText<"Generate SYCL code">;
 def fno_sycl : Flag<["-"], "fno-sycl">, Group<f_Group>, Flags<[NoArgumentUnused, CoreOption]>;
 def fsycl_device_only : Flag<["-"], "fsycl-device-only">, Flags<[CoreOption]>,
   HelpText<"Compile SYCL kernels for device">;

clang/lib/AST/ItaniumMangle.cpp

@@ -2455,6 +2455,67 @@ 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.
@@ -2514,6 +2575,11 @@ 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;
 
@@ -2990,6 +3056,11 @@ 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/lib/AST/QualTypeNames.cpp

@@ -197,7 +197,7 @@ static NestedNameSpecifier *createOuterNNS(const ASTContext &Ctx, const Decl *D,
       // Ignore inline namespace;
       NS = dyn_cast<NamespaceDecl>(NS->getDeclContext());
     }
-    if (NS->getDeclName()) {
+    if (NS && NS->getDeclName()) {
       return createNestedNameSpecifier(Ctx, NS, WithGlobalNsPrefix);
     }
     return nullptr;  // no starting '::', no anonymous

clang/lib/Sema/SemaSYCL.cpp

@@ -1638,6 +1638,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "// This is auto-generated SYCL integration header.\n";
   O << "\n";
 
+  O << "#include <CL/sycl/detail/defines.hpp>\n";
   O << "#include <CL/sycl/detail/kernel_desc.hpp>\n";
 
   O << "\n";
@@ -1651,7 +1652,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   }
   O << "\n";
 
-  O << "namespace cl {\n";
+  O << "__SYCL_INLINE namespace cl {\n";
   O << "namespace sycl {\n";
   O << "namespace detail {\n";
 

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

@@ -0,0 +1,68 @@
+// RUN: %clangxx -fsycl-device-only -fsycl-unnamed-lambda -emit-llvm %s -o %t1.bc
+// RUN: llvm-dis %t1.bc -o - | FileCheck %s
+// RUN: %clangxx -fsycl-device-only -fsycl-unnamed-lambda -emit-llvm %s -DUSE_WRAPPER=1 -o %t2.bc
+// RUN: llvm-dis %t2.bc -o - | FileCheck %s
+
+// Mangling of kernel lambda must be the same for both versions of half
+// CHECK: __unique_stable_name{{.*}} = private unnamed_addr constant [52 x i8] c"_ZTSN2cl4sycl6bufferINS0_4pairIDF16_NS0_5dummyEEEEE\00"
+
+// Helper function to get string returned by __unique_stable_name in LLVM IR
+template <typename T>
+void print() {
+  auto temp = __unique_stable_name(T);
+}
+
+// Helper function to get "print" emitted in device code
+template<typename T, typename F>
+__attribute__((sycl_kernel)) void helper(F f) {
+  print<T>();
+  f();
+}
+
+// Half wrapper, as it defined in SYCL headers
+namespace cl {
+namespace sycl {
+namespace detail {
+namespace half_impl {
+class half {
+public:
+  half operator=(int) {return *this;}
+};
+} // namespace half_impl
+} // namespace detail
+} // namespace sycl
+} // namespace cl
+
+#ifndef USE_WRAPPER
+using half = _Float16;
+#else
+using half = cl::sycl::detail::half_impl::half;
+#endif
+
+// A few more fake data types to complicate the mangling
+namespace cl {
+namespace sycl {
+struct dummy {
+  int a;
+};
+template<typename T1, typename T2>
+struct pair {
+  T1 a;
+  T2 b;
+};
+template <typename T>
+class buffer {
+public:
+  T &operator[](int) const { return value; }
+  mutable T value;
+};
+} // namespace sycl
+} // namespace cl
+
+int main() {
+  cl::sycl::buffer<cl::sycl::pair<half, cl::sycl::dummy>> B1;
+
+  helper<decltype(B1)>([](){});
+
+  return 0;
+}