On Wed, Oct 22, 2014 at 10:34 AM, Thomas Schwinge <tho...@codesourcery.com> wrote: > Hi! > > On Wed, 22 Oct 2014 10:18:49 +0200, Richard Biener > <richard.guent...@gmail.com> 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? > > 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.
Oh, I see. For HSA we simply don't emit an offload variant for code we cannot handle. But only for those parts. So it's only offload or fallback for other devices? Thus also never share work between both for example (run N threads on the CPU and M threads on the offload target)? Richard. >> 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
