[Bug other/71064] New: nvptx offloading: "long double" data type

[tschwinge at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=71064

            Bug ID: 71064
           Summary: nvptx offloading: "long double" data type
           Product: gcc
           Version: 7.0
            Status: UNCONFIRMED
          Keywords: openacc
          Severity: normal
          Priority: P3
         Component: other
          Assignee: unassigned at gcc dot gnu.org
          Reporter: tschwinge at gcc dot gnu.org
                CC: amonakov at gcc dot gnu.org, bernds at gcc dot gnu.org,
                    jakub at gcc dot gnu.org, nathan at gcc dot gnu.org
            Blocks: 70945
  Target Milestone: ---

Offloading to nvptx (that is, OpenACC only, currently) does not support the
"long double" data type:

    int main()
    {
      long double ld;

    #pragma acc parallel copyout(ld)
      ld = 0.0L;

      if (ld != 0.0L)
        __builtin_abort();

      return 0;
    }

... results in: "lto1: fatal error: unsupported mode XF".

The PTX ISA does not support a floating point data type with more than 64-bits.
 (It's generally tuned for speed instead of precision, and does not strive for
full IEEE-754 conformance.)

We can either continue to error out (in a more user-friendly way).  Or,
(perhaps only for "-ffast-math" or similar?) implement "long double" with the
maximum available precision (that is, that of "double"; this will require
conversion at run time of "long double" values at the offloading boundary),
whilst perhaps emitting a warning about the reduced precision.  Unless that's
what's required because GCC generally does it that way, it does not appear
appealing to add a lot of "soft-fp" support code to emulate the target's
higher-precision "long double" data type for the nvptx offloading target.


Referenced Bugs:

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=70945
[Bug 70945] Offloading: compatibility of target and offloading toolchains