replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[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
This commit is contained in:
Alexey Bataev 2019-06-18 20:29:06 +00:00
parent ba43840bfe
commit 9f3a805ee9
2 changed files with 38 additions and 11 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

28
clang/lib/AST/ItaniumMangle.cpp
View File

@ -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;

21
clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp
View File

@ -8,13 +8,15 @@
// CHECK-DAG: [[T:%.+]] = type {{.+}}, fp128,
// CHECK-DAG: [[T1:%.+]] = type {{.+}}, i128, i128,
#ifndef _ARCH_PPC
typedef __float128 BIGTYPE;
#else
typedef long double BIGTYPE;
#endif
struct T {
char a;
#ifndef _ARCH_PPC
__float128 f;
#else
long double f;
#endif
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]]* %