[PATCH] OpenMP: Ensure that offloaded variables are public

[Jakub Jelinek]

On Tue, Nov 30, 2021 at 05:24:49PM +0100, Jakub Jelinek via Gcc-patches wrote:
> Consider in one TU
> 
> static int a = 5;
> static int baz (void) { static int b;
> #pragma omp declare target to (b)
> return ++b; }
> int foo (void) { return ++a + baz (); }
> #pragma omp declare target to (a, foo)
> 
> and
> 
> static int a = 5;
> static int baz (void) { static int b;
> #pragma omp declare target to (b)
> return ++b; }
> int bar (void) { return ++a + baz (); }
> #pragma omp declare target to (a, bar)
> 
> int
> main ()
> {
>   int v;
>   #pragma omp target (from: v)
>   v = foo () + bar ();
> }
> 
> in another one.  This has
> 	.quad	a
> 	.quad	4
> 	.quad	b.0
> 	.quad	4
> in .offload_var_table.  I'd guess this must fail to link or load
> with GCN if it makes them forcibly TREE_PUBLIC.
> 
> Why does the GCN plugin or runtime need to know those vars?
> It needs to know the single array that contains their addresses of course...

Actually, you've done it in ACCEL_COMPILER only, so
I assume linking the above two sources with -fopenmp into a single
binary or shared library will still work because LTO when reading
the byte-code in will remangle the names of those variables to something
where they are unique in that single *.s (or *.ptx) it emits.
But, if you put one of those TUs into a shared library and the other
into another shared library, I don't see how it can work anymore,
because both those ELF objects which will be in data sections of those
libraries might have clashing names.

If GCN can't support static variables (but isn't it ELF?) and there is no
other way than sacrifice offloading from multiple shared libraries or binary
in the same process, it at least shouldn't be done for targets which don't
need it (e.g. PTX) and shouldn't be done in the pass you've done it in
(because that means it will walk all the vars for each function it
processes, rather than just once).  So, better place would be e.g.
offload_handle_link_vars in lto/*.c or so.

	Jakub