This is the mail archive of the
gcc@gcc.gnu.org
mailing list for the GCC project.
Re: LTO remapping/deduction of machine modes of types/decls
On Mon, Jan 02, 2017 at 10:38:54PM +0300, Alexander Monakov wrote:
> On Mon, 2 Jan 2017, Jakub Jelinek wrote:
> > On Mon, Jan 02, 2017 at 09:49:55PM +0300, Alexander Monakov wrote:
> > > On Mon, 2 Jan 2017, Jakub Jelinek wrote:
> > > > If the host has long double the same as double, sure, PTX can use its native
> > > > DFmode even for long double. But otherwise, the storage must be
> > > > transferable between accelerator and host.
> > >
> > > Hm, sorry, the 'must' is not obvious to me: is it known that the OpenMP ARB
> > > would find only this implementation behavior acceptable?
> >
> > long double is not non-mappable type in the spec, so it is supposed to work.
> > The implementation may choose not to offload whenever it sees long
> > double/__float128/_Float128/_Float128x etc.
>
> But this is not something the implementation can properly enforce; consider
>
> long double v;
> char buf[sizeof v];
> #pragma omp target map(from:buf)
> sscanf ("1.0", "%Lf", buf);
> memcpy(&v, buf, sizeof v);
>
> The offloading compiler wouldn't see a 'long double' anywhere, it gets brought
> in at linking stage. So the implementation would have to tag individual
> translation units and see only in the end of linking if the offloaded image
> touches a long double datum anywhere. And as the example shows, it would prevent
> using printf-like functions.
Well, it can, but of course it can be more or much more work. The case where
long double appears in code seen by the offloading compiler is likely more common
and easier to deal with, it is just a matter of handling the various modes.
For the newlib library routines that deal with long double, it can be solved
e.g. by multilibing across the long double choices (but, we have usually
XFmode on x86_64-linux, on powerpc64le-linux it can be the IBM double double
(IFmode) or IEEE quad (KFmode), ...), or by just multilibing the affected
subset of functions or passing info on what long double is some other way
to the printf*/scanf* family of functions.
Jakub