diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -1892,8 +1892,10 @@
     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
--- /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;
+}