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
