On Fri, Aug 23, 2013 at 02:55:27PM +0400, Ilya Verbin wrote: > I'm trying to implement the approach with modified lto-wrapper. > Suppose we have a bytecode of the routine foo, streamed during ompexp pass > into some section, say .gnu.omptarget_foo. > In function lto.c:do_whole_program_analysis() an extra partition should be > created, that will contain bytecode from .gnu.omptarget_foo, right? > As far as I understood, in addition to the bytecode of foo, we should also > stream extra symtab_nodes, and read them somewhere in > lto-cgraph.c:input_symtab(). > This means we should maintain 2 symtabs inside WPA stage - original for host > and new for target?
I don't think we should stream into more than one target section. There should be just .gnu.target_lto section (or whatever other suitable name) and should stream into it: 1) all functions and variables with "omp declare target" attribute 2) the outlined bodies of #pragma omp target turned into *.ompfn functions 3) all the types, symtab etc. needed for that If compiling with -flto, you'll also get everything from the CU streamed into the normal LTO section, otherwise you'll get assembly for the host variables/functions/etc. Then the question is what the plugin should perform with these sections, whether it will compile each input .gnu.target_lto section hunk separately (as in non-LTO mode), or with -flto also LTO them together. Jakub