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
- From: Alexander Monakov <amonakov at ispras dot ru>
- To: Jakub Jelinek <jakub at redhat dot com>
- Cc: gcc at gcc dot gnu dot org, Richard Biener <rguenther at suse dot de>, Vladislav Ivanishin <vlad at ispras dot ru>
- Date: Mon, 2 Jan 2017 22:38:54 +0300 (MSK)
- Subject: Re: LTO remapping/deduction of machine modes of types/decls
- Authentication-results: sourceware.org; auth=none
- References: <alpine.LNX.2.20.13.1612302022360.2185@monopod.intra.ispras.ru> <20170102101931.GB21933@tucnak> <alpine.LNX.2.20.13.1701021755140.2523@monopod.intra.ispras.ru> <20170102155403.GE21933@tucnak> <alpine.LNX.2.20.13.1701022142000.2523@monopod.intra.ispras.ru> <20170102190255.GG21933@tucnak>
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.
Alexander