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

[Alexey Bataev via cfe-commits]

Author: abataev
Date: Tue Jun 18 13:29:06 2019
New Revision: 363734

URL: http://llvm.org/viewvc/llvm-project?rev=363734&view=rev
Log:
[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.

Modified:
    cfe/trunk/lib/AST/ItaniumMangle.cpp
    cfe/trunk/test/OpenMP/nvptx_unsupported_type_codegen.cpp

Modified: cfe/trunk/lib/AST/ItaniumMangle.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ItaniumMangle.cpp?rev=363734&r1=363733&r2=363734&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ItaniumMangle.cpp (original)
+++ cfe/trunk/lib/AST/ItaniumMangle.cpp Tue Jun 18 13:29:06 2019
@@ -2607,17 +2607,33 @@ void CXXNameMangler::mangleType(const Bu
   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;

Modified: cfe/trunk/test/OpenMP/nvptx_unsupported_type_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_unsupported_type_codegen.cpp?rev=363734&r1=363733&r2=363734&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_unsupported_type_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/nvptx_unsupported_type_codegen.cpp Tue Jun 18 13:29:06 2019
@@ -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]]* %