public inbox for gcc@gcc.gnu.org
 help / color / mirror / Atom feed
From: Alexander Monakov <amonakov@ispras.ru>
To: Jakub Jelinek <jakub@redhat.com>
Cc: gcc@gcc.gnu.org, Richard Biener <rguenther@suse.de>,
	    Vladislav Ivanishin <vlad@ispras.ru>
Subject: Re: LTO remapping/deduction of machine modes of types/decls
Date: Mon, 02 Jan 2017 19:39:00 -0000	[thread overview]
Message-ID: <alpine.LNX.2.20.13.1701022216530.2523@monopod.intra.ispras.ru> (raw)
In-Reply-To: <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

  reply	other threads:[~2017-01-02 19:39 UTC|newest]

Thread overview: 16+ messages / expand[flat|nested]  mbox.gz  Atom feed  top
2016-12-30 18:37 Alexander Monakov
2017-01-02 10:19 ` Jakub Jelinek
2017-01-02 15:39   ` Alexander Monakov
2017-01-02 15:54     ` Jakub Jelinek
2017-01-02 18:50       ` Alexander Monakov
2017-01-02 19:03         ` Jakub Jelinek
2017-01-02 19:39           ` Alexander Monakov [this message]
2017-01-02 20:24             ` Jakub Jelinek
2017-01-04 10:14           ` Richard Biener
2017-01-04 10:04   ` Richard Biener
2017-01-09 18:56     ` Alexander Monakov
2017-01-10  9:23       ` Richard Biener
2017-01-10 15:30         ` Alexander Monakov
2017-01-11  8:16           ` Richard Biener
2017-01-11 14:06             ` Alexander Monakov
2017-01-10 15:53         ` Vladislav Ivanishin

Reply instructions:

You may reply publicly to this message via plain-text email
using any one of the following methods:

* Save the following mbox file, import it into your mail client,
  and reply-to-all from there: mbox

  Avoid top-posting and favor interleaved quoting:
  https://en.wikipedia.org/wiki/Posting_style#Interleaved_style

* Reply using the --to, --cc, and --in-reply-to
  switches of git-send-email(1):

  git send-email \
    --in-reply-to=alpine.LNX.2.20.13.1701022216530.2523@monopod.intra.ispras.ru \
    --to=amonakov@ispras.ru \
    --cc=gcc@gcc.gnu.org \
    --cc=jakub@redhat.com \
    --cc=rguenther@suse.de \
    --cc=vlad@ispras.ru \
    /path/to/YOUR_REPLY

  https://kernel.org/pub/software/scm/git/docs/git-send-email.html

* If your mail client supports setting the In-Reply-To header
  via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).