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
