On Wed, Oct 22, 2014 at 10:18:49AM +0200, Richard Biener wrote:
> On Tue, Oct 21, 2014 at 11:32 PM, Bernd Schmidt <ber...@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?
>
> 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.
I don't think we want to do a fine-grained granularity here, it will only
lead to significant nightmares. E.g. a target region can call other target
functions, if a target function it calls (perhaps directly through a series
of other target functions, perhaps indirectly through function pointers
etc.) can't be supported by the host, you'd need to give up on offloading
all target regions that do or could invoke that. That can be in another TU
within the same shared library etc. And, if some regions are emitted and
others are not, #pragma omp target data will behave less predictably and
more confusingly, right now it can test, does this library have usable
offloading for everything it provides (i.e. libgomp would ask the plugin to
initialize offloading from the current shared library if not already done,
and if successful, say it supports offloading for the particular device and
map variables to that device as requested, otherwise it would just assume
only host fallback is possible and not really map anything). When a target
region is hit, from either within the target data region or elsewhere, it is
already figured out if it has to fallback to host or not.
Now, if you have fine-grained offloading, 33.2% of target regions being
offloadable, the rest not, what would you actually do in target data region?
It doesn't generically know what target regions will be encountered.
So act as if offloading perhaps was possible? But then at each target
region find out if it is really possible?
IMHO people that care about performance will use target regions with care,
with the offloading targets that they care about in mind, for those that
don't care about that, either they will be lucky and things will work out
all, or they will just end up with host fallback.
Jakub
