This is the mail archive of the
gcc@gcc.gnu.org
mailing list for the GCC project.
Re: Questions about LTO infrastructure and pragma omp target
- From: Jakub Jelinek <jakub at redhat dot com>
- To: "Michael V. Zolotukhin" <michael dot v dot zolotukhin at gmail dot com>
- Cc: Thomas Schwinge <thomas at codesourcery dot com>, Ilya Verbin <iverbin at gmail dot com>, Richard Biener <richard dot guenther at gmail dot com>, Uday Khedker <uday at cse dot iitb dot ac dot in>, hubicka at ucw dot cz, rth at redhat dot com, kirill dot yukhin at gmail dot com, gcc at gcc dot gnu dot org
- Date: Tue, 3 Sep 2013 20:54:02 +0200
- Subject: Re: Questions about LTO infrastructure and pragma omp target
- Authentication-results: sourceware.org; auth=none
- References: <20130823105527 dot GA6976 at msticlxl57 dot ims dot intel dot com> <85e37f42-69fe-4bbf-bf1d-f73194e7c444 at email dot android dot com> <20130823123638 dot GL1814 at tucnak dot redhat dot com> <20130823171514 dot GB6976 at msticlxl57 dot ims dot intel dot com> <20130823172347 dot GP1814 at tucnak dot redhat dot com> <20130903135935 dot GC43295 at msticlxl57 dot ims dot intel dot com> <20130903141837 dot GF21876 at tucnak dot zalov dot cz> <20130903151801 dot GE43295 at msticlxl57 dot ims dot intel dot com> <87ioyhhkk9 dot fsf at kepler dot schwinge dot homeip dot net> <20130903182956 dot GH43295 at msticlxl57 dot ims dot intel dot com>
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
On Tue, Sep 03, 2013 at 10:29:56PM +0400, Michael V. Zolotukhin wrote:
> > The idea, as we discussed it at the GNU Tools Cauldron's Acceleration
> > BoF, is that the host program (for at least some acceleration devices)
> > will be responsible for loading the acceleration device's code to the
> > device, using some support library that is specific to each acceleration
> > device
> Unfortunately, I missed the Cauldron, though I'm familiar with the
> general idea and now I'm trying to clarify details.
>
> > and for that it is useful to have the the code readily accessible
> > in the host program, and thus link it in as "data".
> Oh, if we just link the target binary as a data section into the host
> binary, then I see no problems in that, it seems absolutely feasible
> with the existing infrastructure. I just thought (seemingly it was
> incorrect) that we're speaking about linking of target code with the
> host code.
No. The rough idea is that you emit the accelerator related subset of CUs
into the (special named) LTO sections, and when linking a binary you collect
all those sections from all the input object files, compile those (without
-flto ideally separately), link together and finally embed into the
executable into a data section. Similarly when linking a shared library,
you do a target shared library and embed it in a data section of the host
shared library. It is kind of fat binaries/shared libraries. Each of the
accelerators would use different name of the data sections, so they could
coexist.
For the MIC, you'd then use COI to create the binary or shared libraries
from the (ro)data section memory image. For others whatever they support.
Perhaps it should be for MIC a shared library even in the binary and have
some binary in a data section of the libgomp plugin, because it really
should work also if the host binary doesn't have any #pragma omp target
in it at all, but shared libraries do.
Jakub