Re: The nvptx port [1/11+] indirect jumps

[Thomas Schwinge <thomas at codesourcery dot com>]

Hi!

On Wed, 22 Oct 2014 10:18:49 +0200, Richard Biener <richard.guenther@gmail.com> wrote:
> On Tue, Oct 21, 2014 at 11:32 PM, Bernd Schmidt <bernds@codesourcery.com> wrote:
> > On 10/21/2014 11:30 PM, Jakub Jelinek wrote:
> >>
> >> At least for OpenMP, the best would be if the #pragma omp target regions
> >> and/or #pragma omp declare target functions contain anything a particular
> >> offloading accelerator can't handle, instead of failing the whole
> >> compilation perhaps just emit some at least by default non-fatal warning
> >> and not emit anything for the particular offloading target, which would
> >> mean
> >> either host fallback, or, if some other offloading target succeeded, just
> >> that target.
> >
> >
> > I guess a test could be added to mkoffload if gcc were to return a different
> > value for a sorry vs. any other compilation failure. The tool could then
> > choose not to produce offloading support for that target.
> 
> But that would be for the whole file instead of for the specific region?

I'm not sure that's what you're suggesting, but at least on non-shared
memory offloading devices, you can't switch arbitrarily between
offloading device(s) and host-fallback, for you have to do data
management between the non-shared memories.

> So maybe we should produce one LTO offload object for each offload
> function and make the symbols they are supposed to provide weak
> so a fail doesn't end up failing to link the main program?
> 
> Looks like this gets somewhat awkward with the LTO setup.


GrÃÃe,
 Thomas

Attachment: pgpz6osBE9DE_.pgp
Description: PGP signature