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

[Ilya Verbin]

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