[OPENMP]Use host's mangling for 128 bit float types on the device.

Device have to use the same mangling as the host for 128bit float types. Otherwise, the codegen for the device is unable to find the parent function when it tries to generate the outlined function for the target region and it leads to incorrect compilation and crash at the runtime.

llvm-svn: 363734

diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp
index 4fe3bfe65615a6212e879dbffd8dc5f5b00853f9..031aac0dbee9b5a78f0f1151cd6389a3a13e708f 100644 (file)
--- a/clang/lib/AST/ItaniumMangle.cpp
+++ b/clang/lib/AST/ItaniumMangle.cpp
@@ -2607,17 +2607,33 @@ void CXXNameMangler::mangleType(const BuiltinType *T) {
   case BuiltinType::Double:
     Out << 'd';
     break;
-  case BuiltinType::LongDouble:
-    Out << (getASTContext().getTargetInfo().useFloat128ManglingForLongDouble()
-                ? 'g'
-                : 'e');
+  case BuiltinType::LongDouble: {
+    bool UseFloat128Mangling =
+        getASTContext().getTargetInfo().useFloat128ManglingForLongDouble();
+    if (getASTContext().getLangOpts().OpenMP &&
+        getASTContext().getLangOpts().OpenMPIsDevice) {
+      UseFloat128Mangling = getASTContext()
+                                .getAuxTargetInfo()
+                                ->useFloat128ManglingForLongDouble();
+    }
+    Out << (UseFloat128Mangling ? 'g' : 'e');
     break;
-  case BuiltinType::Float128:
-    if (getASTContext().getTargetInfo().useFloat128ManglingForLongDouble())
+  }
+  case BuiltinType::Float128: {
+    bool UseFloat128Mangling =
+        getASTContext().getTargetInfo().useFloat128ManglingForLongDouble();
+    if (getASTContext().getLangOpts().OpenMP &&
+        getASTContext().getLangOpts().OpenMPIsDevice) {
+      UseFloat128Mangling = getASTContext()
+                                .getAuxTargetInfo()
+                                ->useFloat128ManglingForLongDouble();
+    }
+    if (UseFloat128Mangling)
       Out << "U10__float128"; // Match the GCC mangling
     else
       Out << 'g';
     break;
+  }
   case BuiltinType::NullPtr:
     Out << "Dn";
     break;

diff --git a/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp b/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp
index 2573daaa4184200f3009886fbd46ec1bfcd2a7a5..81e7df7e35af2cfedac9e7355989610b5abf3fb8 100644 (file)
--- a/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp
@@ -8,13 +8,15 @@
 // CHECK-DAG: [[T:%.+]] = type {{.+}}, fp128,
 // CHECK-DAG: [[T1:%.+]] = type {{.+}}, i128, i128,
 
-struct T {
-  char a;
 #ifndef _ARCH_PPC
-  __float128 f;
+typedef __float128 BIGTYPE;
 #else
-  long double f;
+typedef long double BIGTYPE;
 #endif
+
+struct T {
+  char a;
+  BIGTYPE f;
   char c;
   T() : a(12), f(15) {}
   T &operator+(T &b) { f += b.a; return *this;}
@@ -68,3 +70,12 @@ void baz1() {
   T1 t = bar1();
 }
 #pragma omp end declare target
+
+BIGTYPE foo(BIGTYPE f) {
+#pragma omp target map(f)
+  f = 1;
+  return f;
+}
+
+// CHECK: define weak void @__omp_offloading_{{.+}}foo{{.+}}_l75([[BIGTYPE:.+]]*
+// CHECK: store [[BIGTYPE]] 0xL00000000000000003FFF000000000000, [[BIGTYPE]]* %