[clang] 0195f86 - [Clang] Fix long double availability check

[Qiu Chaofan via cfe-commits]

Author: Qiu Chaofan
Date: 2021-09-17T15:24:06+08:00
New Revision: 0195f8621f1814967f9cd3ef51ee61117e914299

URL: https://github.com/llvm/llvm-project/commit/0195f8621f1814967f9cd3ef51ee61117e914299
DIFF: https://github.com/llvm/llvm-project/commit/0195f8621f1814967f9cd3ef51ee61117e914299.diff

LOG: [Clang] Fix long double availability check

fae0dfa changed code to check 128-bit float availability, since it
introduced a new 128-bit double type on PowerPC. However, there're other
long float types besides IEEE float128 and PPC double-double requiring
this feature.

Reviewed By: ronlieb

Differential Revision: https://reviews.llvm.org/D109943

Added: 
    clang/test/OpenMP/amdgcn_ldbl_check.cpp

Modified: 
    clang/lib/Sema/Sema.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index a9867697a4c31..0b936d60fc5ee 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -1892,8 +1892,10 @@ void Sema::checkDeviceDecl(ValueDecl *D, SourceLocation Loc) {
     bool LongDoubleMismatched = false;
     if (Ty->isRealFloatingType() && Context.getTypeSize(Ty) == 128) {
       const llvm::fltSemantics &Sem = Context.getFloatTypeSemantics(Ty);
-      if (!Ty->isIbm128Type() && !Ty->isFloat128Type() &&
-          &Sem != &Context.getTargetInfo().getLongDoubleFormat())
+      if ((&Sem != &llvm::APFloat::PPCDoubleDouble() &&
+           !Context.getTargetInfo().hasFloat128Type()) ||
+          (&Sem == &llvm::APFloat::PPCDoubleDouble() &&
+           !Context.getTargetInfo().hasIbm128Type()))
         LongDoubleMismatched = true;
     }
 

diff  --git a/clang/test/OpenMP/amdgcn_ldbl_check.cpp b/clang/test/OpenMP/amdgcn_ldbl_check.cpp
new file mode 100644
index 0000000000000..8971da779c12a
--- /dev/null
+++ b/clang/test/OpenMP/amdgcn_ldbl_check.cpp
@@ -0,0 +1,27 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple x86_64-mingw64 -emit-llvm-bc -target-cpu x86-64 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -o %t.bc -x c++ %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -aux-triple x86_64-mingw64 -fsyntax-only -target-cpu gfx900 -fopenmp -fopenmp-is-device -fopenmp-host-ir-file-path %t.bc -x c++ %s
+// expected-no-diagnostics
+
+void print(double);
+
+constexpr double operator"" _X (long double a)
+{
+	return (double)a;
+}
+
+int main()
+{
+	auto a = 1._X;
+  print(a);
+#pragma omp target map(tofrom: a)
+	{
+#pragma omp teams num_teams(1) thread_limit(4)
+		{
+			a += 1._X;
+		}
+	}
+  print(a);
+	return 0;
+}