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

[Richard Biener <richard dot guenther at gmail dot com>]

On Thu, Dec 4, 2014 at 8:35 PM, Ilya Verbin <iverbin@gmail.com> wrote:
> 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

So where is that "magic" target use then?  Why doesn't the symtab
reachability code see the use?  (why don't we prune it from the offload
table?)

Richard.

> 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