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]]* %
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
