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
