[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

1594

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

1594

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

1594

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

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1594

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

1594

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

1594

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

1594

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

1594

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

1594

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

1594

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

1594

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

1594

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

1594

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

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1594

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

1594

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

1594

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 ↗

(On Diff #208286)

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

1599

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

1599

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

1599

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

1599

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

1599

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

1599

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

1599

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

1599

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

1599

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

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1599

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

1599

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

[jdenny]

clang/lib/Sema/SemaOpenMP.cpp

1599

Sounds good. Thanks again.

clang/lib/Sema/SemaOpenMP.cpp

 
 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()) ||
+       !Context.getTargetInfo().hasFloat128Type() &&
+       Context.getTargetInfo().getLongDoubleWidth() != 128) ||
       (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 &&
        !Context.getTargetInfo().hasInt128Type()))
     targetDiag(E->getExprLoc(), diag::err_type_unsupported)
         << 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;
 
   // Find the directive that is associated with the provided scope.

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 powerpc64le-unknown-linux-gnu -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