Re: [PATCH 3/n] OpenMP 4.0 offloading infrastructure: offload tables

[Ilya Verbin]

Hi,

On 30 Sep 18:53, Ilya Verbin wrote:
> This patch creates 2 vectors with decls: offload_funcs and offload_vars.
> libgomp will use addresses from these arrays to look up offloaded code.
> 
> During the compilation they are outputted to:
> * binary __gnu_offload_funcs/vars sections, or using
>   targetm.record_offload_symbol hook for PTX.

In some cases LTO may optimize out a global variable, declared as target, but
it still will be referenced from the offload table, that will cause a linking
error.  Here is the example:

#pragma omp declare target
int G;
#pragma omp end declare target

int main ()
{
  int res = 0;

  #pragma omp target map(alloc: G) map(from: res)
    {
      G = 1;
      res = G;
    }

  return res;
}

$ gcc -fopenmp -flto -O1 test.c
xxx.ltrans0.ltrans.o:.offload_var_table.3973: error: undefined reference to 'G'

This issue can be resolved by forcing output of such variables.
Is this fix ok?  Should I add a testcase?


diff --git a/gcc/varpool.c b/gcc/varpool.c
index 0526b7f..db28c2a 100644
--- a/gcc/varpool.c
+++ b/gcc/varpool.c
@@ -175,6 +175,7 @@ varpool_node::get_create (tree decl)
       g->have_offload = true;
       if (!in_lto_p)
        vec_safe_push (offload_vars, decl);
+      node->force_output = 1;
 #endif
     }
 
 
Thanks,
  -- Ilya