[gomp4] Tweak GOMP_target{,_data,_update} arguments

Hi!

As discussed earlier, I'd like to pass __OPENMP_TARGET__ argument to
all of GOMP_target{,_data,_update}, so that all those functions
can get at the offloading data section in the shared library or binary
making the call, so that the first time they encounter such a call
in the shared library or binary, it can attempt to actually offload
it to the target (and, if that fails, remember it, and do host fallback).
The reason to pass this also to GOMP_target_data and GOMP_target_update is
that even for those calls the global vars will already need to be mapped,
and to make sure e.g. GOMP_target_data doesn't perform on the target device,
while GOMP_target with the same device clause is performed as host fallback
(that could result in wrong code, e.g. when target data copies data from
device to host at the end, while target construct would modify the host
copies, thus the current host copy would be overwritten by stale target
copy).

As weak undefined hidden symbols don't seem to work as I hoped they would
do, I'm just passing NULL for now, we'll need to wait until we have a linker
plugin that will create the offloading section for us and add hidden
__OPENMP_TARGET__ symbol.

2013-09-18  Jakub Jelinek  <jakub@redhat.com>

	* omp-builtins.def (BUILT_IN_GOMP_TARGET_DATA,
	BUILT_IN_GOMP_TARGET_UPDATE): Use BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR
	rather than BT_FN_VOID_INT_SIZE_PTR_PTR_PTR.
	* builtin-types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR): Remove.
	(BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR): New.
	* omp-low.c (expand_omp_target): Remove fnname argument from
	GOMP_target, add openmp_target argument to GOMP_target,
	GOMP_target_data and GOMP_target_update calls.
gcc/fortran/
	* types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR): Remove.
	(BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR): New.
libgomp/
	* libgomp_g.h (GOMP_target): Change 3rd argument to const void *.
	(GOMP_target_data, GOMP_target_update): Add openmp_target argument.
	* target.c (GOMP_target): Change fnname argument to openmp_target,
	with const void * type.
	(GOMP_target_data, GOMP_target_update): Add openmp_target argument.


	Jakub

> As discussed earlier, I'd like to pass __OPENMP_TARGET__ argument to
> all of GOMP_target{,_data,_update}, so that all those functions
> can get at the offloading data section in the shared library or binary
> making the call, so that the first time they encounter such a call
> in the shared library or binary, it can attempt to actually offload
> it to the target (and, if that fails, remember it, and do host fallback).
Actually, I thought that that we'd perform loading the target code in
gomp-initialization.  So, first time libgomp is called, it loads target image
(from somewhere) and starts the target-side MAIN on the target (which actually
just starts to wait for offload requests).  Then, when GOMP_target is called, it
just passes address (or name, as it was initially) to this target-side MAIN and
the target executes the requested function.  I thought that at the moment
GOMP_target is called everything is ready and loaded as initialization has been
already executed.  Isn't that correct?  Maybe I misunderstood something - do you
want to load target images only in GOMP_target, not earlier?

And I still not quite get it, how would target code know, which function to call
in this scheme.  How will it figure out, which function to call if we didn't
even pass the function name?  As far as I understand, __OPENMP_TARGET__ will
lead to a dedicated section, which contains addresses of host-side versions of
the functions.  But how could we obtain the corresponding addresses for the
target?  COI, e.g., as I understand it, currently doesn't have a capability of
returning target-side function address - it only could call a function by name.
Or am I missing something?

Thanks, Michael
> 	Jakub

On Wed, Sep 18, 2013 at 06:13:25PM +0400, Michael V. Zolotukhin wrote:
> > As discussed earlier, I'd like to pass __OPENMP_TARGET__ argument to
> > all of GOMP_target{,_data,_update}, so that all those functions
> > can get at the offloading data section in the shared library or binary
> > making the call, so that the first time they encounter such a call
> > in the shared library or binary, it can attempt to actually offload
> > it to the target (and, if that fails, remember it, and do host fallback).
> Actually, I thought that that we'd perform loading the target code in
> gomp-initialization.  So, first time libgomp is called, it loads target image
> (from somewhere) and starts the target-side MAIN on the target (which actually
> just starts to wait for offload requests).  Then, when GOMP_target is called, it
> just passes address (or name, as it was initially) to this target-side MAIN and
> the target executes the requested function.  I thought that at the moment
> GOMP_target is called everything is ready and loaded as initialization has been
> already executed.  Isn't that correct?  Maybe I misunderstood something - do you
> want to load target images only in GOMP_target, not earlier?

Yes, I want to load target images only in GOMP_target{,_data,_update}.
Loading it earlier is both premature (if you have say two MICs, one HSAIL
available, why would you upload everything to all 3 devices, even when you
don't know if the program will just print usage and exit and never enter any
target region), and hard to do (a program can have many shared libraries
(plus the binary), and each of them can have their offload stuff, how do you
find out what to offload and what not, and from where?).

__OPENMP_TARGET__ would be a linker plugin inserted symbol at the start of
some linker plugin created data section, which would start with some header
and then data.
Say
uleb128 number_of_supported_targets - n
uleb128 number_of_host_var_pairs - m
[ name of offload target (asciiz?)
  relative offset to the start of the offload data for the target (in MIC case embedded DSO)
  size of the offload data
  perhaps something how to find the target addresses array
] repeated n times
[ host_address, size ] repeated m times
(for the functions passed to GOMP_target the pair would be [ foobar.omp_fn.25, 1 ] ).

If no offload data is provided at all, then the section would be just one
byte - 0 (or two bytes, 0 0)?

So, when GOMP_target{,_data,_update} is called, it could easily determine
if the calling shared library resp. binary has been offloaded or not (it
could e.g. have a few entries array as cache, otherwise lookup the
__OPENMP_TARGET__ address in the splay tree or something?
If it sees it hasn't been initialized yet, if it is the very first time,
for MIC it would upload the binary (from the plugin data section?), libgomp
and finally the shared library from the __OPENMP_TARGET__ and initialize the
splay tree with the variables.

> And I still not quite get it, how would target code know, which function to call
> in this scheme.  How will it figure out, which function to call if we didn't
> even pass the function name?  As far as I understand, __OPENMP_TARGET__ will

See earlier mail, function names aren't neccessarily unique.  After the
__OPENMP_TARGET__ shared library is dlopened in the target and splay tree
initialized from the host array of [host_address, size] pairs and
corresponding array of target_address in the target code, you just lookup
the [ fnaddr, 1 ] in the splay tree and use the target address from that
mapping.

> lead to a dedicated section, which contains addresses of host-side versions of
> the functions.  But how could we obtain the corresponding addresses for the
> target?  COI, e.g., as I understand it, currently doesn't have a capability of
> returning target-side function address - it only could call a function by name.
> Or am I missing something?

	Jakub

Hi Jakub,

Thanks for the explanation, it's getting a bit clearer, though I still have some
questions.

> __OPENMP_TARGET__ would be a linker plugin inserted symbol at the start of
> some linker plugin created data section, which would start with some header
> and then data.
> Say
> uleb128 number_of_supported_targets - n
> uleb128 number_of_host_var_pairs - m
> [ name of offload target (asciiz?)
>   relative offset to the start of the offload data for the target (in MIC case embedded DSO)
>   size of the offload data
>   perhaps something how to find the target addresses array
> ] repeated n times
> [ host_address, size ] repeated m times
> (for the functions passed to GOMP_target the pair would be [ foobar.omp_fn.25, 1 ] ).
So, in this table we store host addresses of global variables, marked with
'pragma omp declare target', and addresses of host-versions of OMP-versioned
functions.  Correct?  Also, there are pointers to images of target-binaries,
which are (presumably) placed in other (or the same?) data sections.

> So, when GOMP_target{,_data,_update} is called, it could easily determine
> if the calling shared library resp. binary has been offloaded or not
That's right.
Then, if no initialization has been performed yet, GOMP_target{,_data,_update}
is initialized.  Now let's look at the initialization.
In initialization GOMP_target* looks at the __OPENMP_TARGET__ table (its address
is passed as the 3rd argument), finds pointer to a data section with
target-binary image, loads it to memory, runs a process on a target from it
(e.g. in COI using COIProcessCreateFromFile and/or
COIProcessLoadLibraryFromMemory).
Global variables are mapped and the corresponding host<->target address pairs
are inserted to the splay tree, as usual.
Also, GOMP_target* should do the same for function addresses.  Could you please
describe this step in more details?  Do we want to just add some offset to
host_function_address (as we want host versions of functions to be ordered
exactly as the target versions)?

> See above, names are just a bad idea.  You can just use some magic wrapper
> name in the target binary (the one sitting in libgomp), to which you just
> pass the pair of function address and it's argument and the named function
> will just read the (target) function pointer and (target) pointer argument
> from misc data block and tail call that function.
Yes, if we know target function pointer, we can do this.

Basically, the main question I have now is how would we figure out target
function address?  Of course, after initialization we just look for it in our
splay tree, so the question relates to the initialization step.

Thanks, Michael

On Thu, Sep 19, 2013 at 12:58:28PM +0400, Michael V. Zolotukhin wrote:
> Thanks for the explanation, it's getting a bit clearer, though I still have some
> questions.
> 
> > __OPENMP_TARGET__ would be a linker plugin inserted symbol at the start of
> > some linker plugin created data section, which would start with some header
> > and then data.
> > Say
> > uleb128 number_of_supported_targets - n
> > uleb128 number_of_host_var_pairs - m
> > [ name of offload target (asciiz?)
> >   relative offset to the start of the offload data for the target (in MIC case embedded DSO)
> >   size of the offload data
> >   perhaps something how to find the target addresses array
> > ] repeated n times
> > [ host_address, size ] repeated m times
> > (for the functions passed to GOMP_target the pair would be [ foobar.omp_fn.25, 1 ] ).
> So, in this table we store host addresses of global variables, marked with
> 'pragma omp declare target', and addresses of host-versions of OMP-versioned
> functions.  Correct?  Also, there are pointers to images of target-binaries,
> which are (presumably) placed in other (or the same?) data sections.

Yeah.  How exactly we define the section is up to us, but it should have all
the information that GOMP_target* will need to offload the stuff from the
current shared library or binary, and everything needed to initialize the
{ host_addr, size } -> { target_addr } mapping of declare target global var
definitions and functions passed to GOMP_target.  The fewer relocations
the section has, the better.  But, if we need any relocations, it will need
to be in a relro section, and supposedly the embedded shared library (resp.
libraries) don't need any relocations on them and will be large, thus
supposedly they should live in different sections and the header should just
point to them (e.g. using offset relative to __OPENMP_TARGET__ or something
that doesn't dynamic relocation).  Similarly, if the linker plugin puts in the
array of [ host_address, size ] rewritten such that host_address is an
offset from __OPENMP_TARGET__, then we won't need dynamic relocations for
that.  Another complication is dependent shared libraries.
Consider
liba.c:
#pragma omp declare target
int i;
int foo (void)
{
  return ++i;
}
#pragma omp end declare target
main.c:
#pragma omp declare target
extern int i;
extern int foo (void);
#pragma omp end declare target
int main ()
{
  int j;
  #pragma omp target
    {
      j = i;
      j += foo ();
    }
  if (j != 1)
    abort ();
  return 0;
}
gcc -shared -O2 -fpic -fopenmp -o liba.so -Wl,-soname,liba.so liba.c
gcc -O2 -fopenmp -o main main.c -L. -la
./main

Perhaps the linker plugin can extract the target shared libraries from
the embedded sections of dependent shared libraries (if any), and link the
"main" shared library against that, but GOMP_target will need to know that
it can't just offload main.so, but also has to offload the dependent
liba.so (and of course libgomp.so.1 from the libgomp plugin).
What does ICC do in this case?

> > So, when GOMP_target{,_data,_update} is called, it could easily determine
> > if the calling shared library resp. binary has been offloaded or not
> That's right.
> Then, if no initialization has been performed yet, GOMP_target{,_data,_update}
> is initialized.  Now let's look at the initialization.
> In initialization GOMP_target* looks at the __OPENMP_TARGET__ table (its address
> is passed as the 3rd argument), finds pointer to a data section with
> target-binary image, loads it to memory, runs a process on a target from it
> (e.g. in COI using COIProcessCreateFromFile and/or
> COIProcessLoadLibraryFromMemory).
> Global variables are mapped and the corresponding host<->target address pairs
> are inserted to the splay tree, as usual.
> Also, GOMP_target* should do the same for function addresses.  Could you please
> describe this step in more details?  Do we want to just add some offset to
> host_function_address (as we want host versions of functions to be ordered
> exactly as the target versions)?

The idea was that the host [ host_addr, size ] array (in some named section)
would be ordered exactly the same as corresponding [ targ_addr ] array in
the target shared library.  So, [25] pair in the host array will correspond
to [25] in the target shared library array.
So you just walk the whole arrays, and in each iteration pick nth host array
pair plus corresponding nth target array address, and put it into the splay
tree.

In the above testcase, host liba.so would contain a [ &i, sizeof(int) ]
pair and target liba.so corresponding [ &i ] entry (target i in that case).
In host main there would be [ &main.omp_fn.0, 1 ] and in target main.so
corresponding [ &main.omp_fn.0 ] (target main.omp_fn.0 in that case).

> > See above, names are just a bad idea.  You can just use some magic wrapper
> > name in the target binary (the one sitting in libgomp), to which you just
> > pass the pair of function address and it's argument and the named function
> > will just read the (target) function pointer and (target) pointer argument
> > from misc data block and tail call that function.
> Yes, if we know target function pointer, we can do this.
> 
> Basically, the main question I have now is how would we figure out target
> function address?  Of course, after initialization we just look for it in our
> splay tree, so the question relates to the initialization step.

After you ensure the shared library is offloaded and splay tree initialized,
you just splay_tree_lookup the host [ fnaddr, 1 ] and get corresponding
target address (if not present in splay tree, that would be toolchain bug,
so gomp_fatal or something).  And, then you just make sure you call that
function in the target, whether it is done by calling a named wrapper
function to which you pass that target fn address and target address of the
pointer array, or something else.

	Jakub

On 19 Sep 11:23, Jakub Jelinek wrote:
> that.  Another complication is dependent shared libraries.
> Consider
> liba.c:
> #pragma omp declare target
> int i;
> int foo (void)
> {
>   return ++i;
> }
> #pragma omp end declare target
> main.c:
> #pragma omp declare target
> extern int i;
> extern int foo (void);
> #pragma omp end declare target
> int main ()
> {
>   int j;
>   #pragma omp target
>     {
>       j = i;
>       j += foo ();
>     }
>   if (j != 1)
>     abort ();
>   return 0;
> }
> gcc -shared -O2 -fpic -fopenmp -o liba.so -Wl,-soname,liba.so liba.c
> gcc -O2 -fopenmp -o main main.c -L. -la
> ./main
> 
> Perhaps the linker plugin can extract the target shared libraries from
> the embedded sections of dependent shared libraries (if any), and link the
> "main" shared library against that, but GOMP_target will need to know that
> it can't just offload main.so, but also has to offload the dependent
> liba.so (and of course libgomp.so.1 from the libgomp plugin).
> What does ICC do in this case?
> 
> 	Jakub

Hi Jakub,

Here's what ICC does.
Suppose we have liba.c and main.c, both with target regions:

1. Building liba.c -> liba.so.
A call to offload-runtime library is inserted into _init of liba.so.
Target region is compiled into liba_target.so, and placed into .rodata of
liba.so.

2. Building main.c -> main.exe.
Similarly, a call to offload-runtime library is inserted into _init of main.exe.
Target region is compiled into main_target.so, and placed into .rodata of
main.exe.

3. Runtime.
So, when liba.so and main.exe are loaded at host-side, the runtime library
knows, that it should transfer liba_target.so and main_target.so to the
target-side.  Then, main.exe starts execution.  At every entry point to the
target region, runtime library checks whether it should perform an
initialization.  If target is not initialized, runtime library calls
COIProcessCreateFromMemory(main_target.exe), that transfers some standard
main_target.exe to the target and starts it.  Then, runtime library calls
COIProcessLoadLibraryFromMemory(liba_target.so, main_target.so), that transfers
these libraries to the target and loads them into the main_target.exe.
The target-side functions are called from host through
COIProcessGetFunctionHandles("f_name") and COIPipelineRunFunction(handle). The
addresses of target-side functions are obtained from *_target.so by dlsym().
So, the host-side knows nothing about target addresses.

What do you think, how will such an approach work with other target
architectures, and with current implementation of GOMP_target{,_data,_update}?

Thanks,
  -- Ilya