[OpenMP] Fix 128-bit long double support on target

[jdenny]

Summary

For example, without this fix, when the host is x86_64, long double is
sometimes rejected when offloading to x86_64:

$ cat test.c
int main() {
  long double in = 0;
  #pragma omp target
    in *= 2;
  return 0;
}
$ clang -fopenmp test.c; echo $?
0
$ clang -fopenmp -fopenmp-targets=x86_64 test.c
test.c:4:5: error: 'long double' is not supported on this target
    in *= 2;
    ^~
1 error generated.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

1592

Hmm, this look strange, at least. Seems to me, in this case the size of the long double is 128 bit (copied from the host), but device reports that it does not support 128 bit double. Seems to me, it is a problem with the device configuration. Why does the host translate long double to 128 bit fp, while the device translates it to 64 bit FP?

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1592

Sorry, I think I've misunderstood what's happening here, and my fix is probably wrong.

For x86_64, the example from my patch summary fails as described there. Does that work for you?

For powerpc64le, the reproducer I added to the test suite fails without this patch. Shouldn't it succeed?

clang/test/OpenMP/nvptx_unsupported_type_messages.cpp

9

Sorry, this line was supposed to be x86_64 only, and then it no longer acts as a reproducer for me.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

1592

Still, seems to me like the problem with the device config, not the original check.

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1592

Still, seems to me like the problem with the device config, not the original check.

I'm not sure where to begin looking for that. Can you point me in the right direction? Thanks.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

1592

You need to understand why host and device report different size of the type. Check how the device is configured in lib/Basic/Targets

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1592

Thanks for the pointer. I think I understand things a bit better now.

Without this patch's fix, the x86_64 example from this patch's summary fails while this patch's new x86_64 test case passes. The difference is the summary's example doesn't specify -unknown-linux after x86_64, and that's what sets hasFloat128Type() to true.

powerpc64le-unknown-linux-gnu does not have __float128, it seems. That's why this patch's new powerpc64le test case fails without this patch's fix.

It seems strange to me that the code we're commenting on originally looks for the source type to be either __float128 or 128-bit long double, and it then requires the target to support __float128. It doesn't accept 128-bit long double support as sufficient. My intention in this patch was to extend it to accept either so that all the examples above compile. Is that too lenient? Am I misinterpreting what's happening?

As for your comment about 64-bit floating point in the device translation, I haven't seen that yet. Did I miss it?

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

1592

The intention of the original patch is to make host and device to have the same float128 and long double types. Device inherits those types from the host to be compatible during offloading and to correctly mangle functions.
Without this we just can't generate offloading regions correctly. If the host has 128 bit long double, the device also must have 128 bit long double.
If device does not support 128bit floats, in this case device can only move the data (do load/stores ops only) and cannot do anything else.

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1592

Are you intentionally requiring support for __float128 when the source type is 128-bit long double? That seems to mean powerpc64le cannot offload to itself.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

1592

No, if the host has 128 bit long double, the device must also have 128 bit long double. It has nothing to do with the float128 type itself.

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1592

What if we change the logic to the following?

(Ty->isFloat128Type() && !Context.getTargetInfo().hasFloat128Type()) ||
(!Ty->isFloat128Type() && Ty->isRealFloatingType() &&
 Context.getTypeSize(Ty) == 128 &&
 Context.getTargetInfo().getLongDoubleWidth() != 128)

Maybe there's a more succinct way to check if Ty is long double....

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

1592

What if Ty is not long double, but some other FP type?

[MaskRay]

Drive-by comment :)

In gcc, (1) -mlong-double-128 uses IEEE 754 quad as the representation of long double on x86_64 (2) -mlong-double-128 -mabi=ieeelongdouble uses IEEE 754 quad as the representation of long double on powerpc{32,64}. (On Linux powerpc{32,64}, -mlong-double-128 -mabi=ibmlongdouble is the default, it uses a IBM 128-bit extended precision for long double).

-mlong-double-128 is currently not supported by clang, but I intend to support it in D64277. After that, the representation of long double will be the same as __float128 (unfortunately they will use the same mangling scheme)

As I noticed, the mangling scheme of __float128 also needs a fix. After applying D64277 and its dependent revisions, the last two CHECK lines of test/OpenMP/nvptx_unsupported_type_codegen.cpp will fail...

[MaskRay]

clang/lib/Sema/SemaOpenMP.cpp

1592

I know little about OpenMP... but does these lines take into account of 128-bit IBM extended double on powerpc{32,64}? It is the default representation of long double.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

1592

Yes, it does, it checks for any 128bit FP type.

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1592

What if Ty is not long double, but some other FP type?

We could use something like this to be sure it's really long double:

Ty.getUnqualifiedType() == Context.LongDoubleTy

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

1592

What if we have the problem with the FP type, not the long double? Shall we miss the check for it? I don't think so. What we need to improve, is the error message to describe that there is a problem with 128 bit tyoe on the given platform.

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1592

What if we have the problem with the FP type, not the long double? Shall we miss the check for it? I don't think so.

I see there's a llvm::fltSemantics. Perhaps there's a way to compare that to be sure representations are equivalent.

What we need to improve, is the error message to describe that there is a problem with 128 bit tyoe on the given platform.

I agree the error message could be misleading as written now: it might claim long double isn't supported when really it's just a different long double.

[jdenny]

This incorporates all the improvements I think we've discussed so far:

• In diagnostics, distinguish between unsupported types and non-equivalent types.
• Don't treat __float128 and long double as equivalent.
• Compare exact floating point representation instead of just size.

[ABataev]

clang/include/clang/Basic/DiagnosticSemaKinds.td

8515

No, this message shall tell something like host requires %0 bit size %1 type support, but device %3 does not support it. Here %0 is the sizr of the host type in bits, %1 is the type name and %2 is the device triple.

clang/lib/Sema/SemaOpenMP.cpp

1596

Why do we need all this stuff? I think the original code works good here, we just need to improve the message.

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1596

Why do we need all this stuff? I think the original code works good here, we just need to improve the message.

It seems we've been miscommunicating. Let's review the discussion so far, point by point, and I'll show you how I arrived at this code. I think the first major point is as follows.

I asked:

Are you intentionally requiring support for __float128 when the source type is 128-bit long double? That seems to mean powerpc64le cannot offload to itself.

You replied:

No, if the host has 128 bit long double, the device must also have 128 bit long double. It has nothing to do with the float128 type itself.

I thought you were agreeing with my understanding. That is, the original code requires __float128 support even when 128-bit long double is in use. That's why powerpc64le cannot offload to itself. How does the original code require __float128? It checks Context.getTargetInfo().hasFloat128Type(). As far as I can tell, that checks for __float128 support, and it does not check for 128-bit long double support. That's why powerpc64le cannot offload to itself.

I'll review other points later so we can discuss them. First, let's see if we can agree on this point.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

1596

My point about ppc64le is:
If the host code uses long double 128 bit long, tbe device long double also must be 128 bit long. But if device does not support 128bit FP type naturally, user cannot do any operations with it except just load/stores.
Forget about float128 type, we talk about 128bit long double here.

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1596

This patch adds new testing for powerpc64le. Without the rest of this patch (and without the expected-error change), should that pass? It does not for me.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

1596

Just like I said before, it means that your device config does not match the expected one. It means that either you specified incorrect triple, or the device is not configured properly.

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1596

The triple is in the new test code. Is it incorrect?

When you discussed device configuration earlier, you made a comment about 64 bit FP, and I couldn't find evidence of that anywhere.

When I explore the device config as you suggested earlier, as far as I can tell, the only issue here is that it claims powerpc64le does not support __float128, but the original logic in Sema::checkOpenMPDeviceExpr requires __float128 by checking Context.getTargetInfo().hasFloat128Type().

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

1596

Original ppc64le host reports that it supports 128bit float. The device reports that it does not support it. It just means that the device has different configuration than host. If yoh want the device to support 128 bit FP as the host, you need either to specify the device correctly, orfix the target device configuration in the lib/Basic/Target if it does not match the expected behavior.
Check the host init for ppc64le. It sets hasFloat128 to true. For some reason, device target config does not set this flag and reports that it does not support 128bit FP.

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1596

Original ppc64le host reports that it supports 128bit float.

Check the host init for ppc64le. It sets hasFloat128 to true.

Doesn't appear to happen for me unless I add -target-feature +float128 as a -cc1 option, as some existing powerpc64le clang tests do.

@MaskRay, will powerpc64le set HasFloat128=true after your changes?

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

1596

This is what I meant! We just missed some special target config options for the device.

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1596

Thanks for the discussion. I know what to do for the x86_64 case. Others seem to be working on powerpc64le. Maybe all architectures we care about (will eventually) correctly support __float128. I'll abandon this patch.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

1596

No problems. I'll fix the error message tomorrow to be more specific.

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1596

Sounds good. Thanks again.

clang/include/clang/AST/ASTContext.h

static bool isObjCNSObjectType(QualType Ty) {
     return Ty->isObjCNSObjectType();
   }
 
   //===--------------------------------------------------------------------===//
   //                         Type Sizing and Analysis
   //===--------------------------------------------------------------------===//
 
   /// Return the APFloat 'semantics' for the specified scalar floating
-  /// point type.
-  const llvm::fltSemantics &getFloatTypeSemantics(QualType T) const;
+  /// point type.  Get the semantics from the auxiliary target if one is
+  /// defined and \p PreferAuxTarget.
+  const llvm::fltSemantics &getFloatTypeSemantics(
+      QualType T, bool PreferAuxTarget = false) const;
 
   /// Get the size and alignment of the specified complete type in bits.
   TypeInfo getTypeInfo(const Type *T) const;
   TypeInfo getTypeInfo(QualType T) const { return getTypeInfo(T.getTypePtr()); }
 
   /// Get default simd alignment of the specified complete type in bits.
   unsigned getOpenMPDefaultSimdAlign(QualType T) const;
 

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def ext_c99_array_usage : Extension<
   "%select{qualifier in |static |}0array size %select{||'[*] '}0is a C99 "
   "feature">, InGroup<C99>;
 def err_c99_array_usage_cxx : Error<
   "%select{qualifier in |static |}0array size %select{||'[*] '}0is a C99 "
   "feature, not permitted in C++">;
 def err_type_unsupported : Error<
   "%0 is not supported on this target">;
+def err_type_not_equivalent : Error<
+  "%0 is not equivalent between host and target">;
 def err_nsconsumed_attribute_mismatch : Error<
   "overriding method has mismatched ns_consumed attribute on its"
   " parameter">;
 def err_nsreturns_retained_attribute_mismatch : Error<
   "overriding method has mismatched ns_returns_%select{not_retained|retained}0"
   " attributes">;
 def warn_nsconsumed_attribute_mismatch : Warning<
   err_nsconsumed_attribute_mismatch.Text>, InGroup<NSConsumedMismatch>;

clang/lib/AST/ASTContext.cpp

 }
 
 //===----------------------------------------------------------------------===//
 //                         Type Sizing and Analysis
 //===----------------------------------------------------------------------===//
 
 /// getFloatTypeSemantics - Return the APFloat 'semantics' for the specified
 /// scalar floating point type.
-const llvm::fltSemantics &ASTContext::getFloatTypeSemantics(QualType T) const {
+const llvm::fltSemantics &ASTContext::getFloatTypeSemantics(
+    QualType T, bool PreferAuxTarget) const {
+  const TargetInfo *Tgt = PreferAuxTarget && AuxTarget ? AuxTarget : Target;
   const auto *BT = T->getAs<BuiltinType>();
   assert(BT && "Not a floating point type!");
   switch (BT->getKind()) {
   default: llvm_unreachable("Not a floating point type!");
   case BuiltinType::Float16:
   case BuiltinType::Half:
-    return Target->getHalfFormat();
-  case BuiltinType::Float:      return Target->getFloatFormat();
-  case BuiltinType::Double:     return Target->getDoubleFormat();
-  case BuiltinType::LongDouble: return Target->getLongDoubleFormat();
-  case BuiltinType::Float128:   return Target->getFloat128Format();
+    return Tgt->getHalfFormat();
+  case BuiltinType::Float:      return Tgt->getFloatFormat();
+  case BuiltinType::Double:     return Tgt->getDoubleFormat();
+  case BuiltinType::LongDouble: return Tgt->getLongDoubleFormat();
+  case BuiltinType::Float128:   return Tgt->getFloat128Format();
   }
 }
 
 CharUnits ASTContext::getDeclAlign(const Decl *D, bool ForAlignof) const {
   unsigned Align = Target->getCharWidth();
 
   bool UseAlignAttrOnly = false;
   if (unsigned AlignFromAttr = D->getMaxAlignment()) {

clang/lib/Sema/SemaOpenMP.cpp

else if (Caller)
     DeviceCallGraph[Caller].insert({Callee, Loc});
 }
 
 void Sema::checkOpenMPDeviceExpr(const Expr *E) {
   assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
          "OpenMP device compilation mode is expected.");
   QualType Ty = E->getType();
   if ((Ty->isFloat16Type() && !Context.getTargetInfo().hasFloat16Type()) ||
-      ((Ty->isFloat128Type() ||
-        (Ty->isRealFloatingType() && Context.getTypeSize(Ty) == 128)) &&
-       !Context.getTargetInfo().hasFloat128Type()) ||
+      (Ty->isFloat128Type() && !Context.getTargetInfo().hasFloat128Type()) ||
       (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 &&
        !Context.getTargetInfo().hasInt128Type()))
     targetDiag(E->getExprLoc(), diag::err_type_unsupported)
         << Ty << E->getSourceRange();
+  if (Ty->isRealFloatingType()) {
+    llvm::APFloatBase::Semantics Sem = llvm::APFloatBase::SemanticsToEnum(
+        Context.getFloatTypeSemantics(Ty, /*PreferAuxTarget*/ true));
+    const TargetInfo &TI = Context.getTargetInfo();
+    if ((Ty->isFloat128Type() &&
+         Sem != llvm::APFloatBase::SemanticsToEnum(TI.getFloat128Format())) ||
+        (Ty.getUnqualifiedType() == Context.LongDoubleTy &&
+         Sem != llvm::APFloatBase::SemanticsToEnum(TI.getLongDoubleFormat())))
+      // TODO: Naming the floating point representations would be helpful.
+      targetDiag(E->getExprLoc(), diag::err_type_not_equivalent)
+          << Ty << E->getSourceRange();
+  }
 }
 
 bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level) const {
   assert(LangOpts.OpenMP && "OpenMP is not allowed");
 
   ASTContext &Ctx = getASTContext();
   bool IsByRef = true;
 

clang/test/OpenMP/nvptx_unsupported_type_messages.cpp

 // Test target codegen - host bc file has to be created first.
-// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -verify=quiet -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -fsyntax-only
-// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-linux-gnu -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -verify=quiet -fopenmp -x c++ -triple powerpc64le-unknown-linux-gnu -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-linux-gnu -fopenmp-targets=nvptx64-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -fsyntax-only
+//
+// Make sure this code does compile when the host and target are the same.
+// RUN: %clang_cc1 -verify=quiet -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=x86_64-unknown-linux -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -verify=quiet -fopenmp -x c++ -triple x86_64-unknown-linux -aux-triple x86_64-unknown-linux -fopenmp-targets=x86_64-unknown-linux %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -fsyntax-only
+// RUN: %clang_cc1 -verify=quiet -fopenmp -x c++ -triple powerpc64le-unknown-linux-gnu -fopenmp-targets=powerpc64le-unknown-linux-gnu -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -verify=quiet -fopenmp -x c++ -triple powerpc64le-unknown-linux-gnu -aux-triple powerpc64le-unknown-linux-gnu -fopenmp-targets=powerpc64le-unknown-linux-gnu %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -fsyntax-only
+
+// quiet-no-diagnostics
 
 struct T {
   char a;
 #ifndef _ARCH_PPC
   __float128 f;
 #else
   long double f;
 #endif
   char c;
   T() : a(12), f(15) {}
 #ifndef _ARCH_PPC
 // expected-error@+4 {{'__float128' is not supported on this target}}
 #else
-// expected-error@+2 {{'long double' is not supported on this target}}
+// expected-error@+2 {{'long double' is not equivalent between host and target}}
 #endif
   T &operator+(T &b) { f += b.a; return *this;}
 };
 
 struct T1 {
   char a;
   __int128 f;
   __int128 f1;