[PATCH] OpenMP: Ensure that offloaded variables are public

Jakub Jelinek jakub@redhat.com
Tue Nov 30 16:24:49 GMT 2021


On Tue, Nov 16, 2021 at 11:49:18AM +0000, Andrew Stubbs wrote:
> This patch is needed for AMD GCN offloading when we use the assembler from
> LLVM 13+.
> 
> The GCN runtime (libgomp+ROCm) requires that the location of all variables
> in the offloaded variables table are discoverable at runtime (using the
> "hsa_executable_symbol_get_info" API), and this only works when the symbols
> are exported from the binary. Previously we solved this by having mkoffload
> insert ".global" directives into the assembler text, but newer LLVM
> assemblers emit an error if we do this when then variable was previously
> declared ".local" (which happens when a variable is zero-initialized and
> placed in the BSS).
> 
> Since we can no longer easily fix them up after the fact, this patch fixes
> them up during OMP lowering.

I'm confused, how can that ever work reliably?
The !TREE_PUBLIC offload_vars can be static locals or static globals
or static anon namespace vars, but their names can very easily clash with
either static or non-static variables from other TUs.
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...

	Jakub



More information about the Gcc-patches mailing list