Re: "omp declare target" on DECL_EXTERNAL vars

[Alexander Monakov <amonakov@ispras.ru>]

On Mon, 23 May 2016, Jakub Jelinek wrote:
> Having the externs specified in omp declare target to is important for
> code generation, we need to know that whether the vars should be mapped
> implicitly on target constructs and remapped in the target construct bodies,
> or whether the actual vars should be used in the regions.

Yep, sorry for missing that.

> Thus, 
> 
> > So from that perspective it's undesirable to have 'omp declare target' on
> > declarations that don't define anything.
> 
> is just wrong, we at least need the symbol_table::offloadable bit set.

So unlike for functions, for variables GCC needs to know exactly whether they
are 'omp declare target [link]' at all points of use, not just at the point of
definition.

There's a pitfall if the user forgets the pragma on the external declaration:

=== a.c

#pragma omp declare target
int a;
void set_a()
{
  a = 42;
}
#pragma omp end declare target

=== main.c

extern int a;
extern void set_a();
#pragma omp declare target to(set_a)

int main()
{
  a = 0;
  #pragma omp target map(tofrom:a)
    set_a();

  if (a != 42) abort();
}
===

As I understand, this aborts, and it's not obvious how to take measures to
produce a compile-time diagnostic.  And I'm not sure if the letter of the spec
is being violated there.

Sorry if I'm elaborating on the more obvious stuff without contributing to
your original question; I hope this is of some value (like it is for me).

> About g->head_offload and offload_vars, I guess it is fine not to set those
> for externs but we need to arrange that to be set when we actually define it
> when it has been previously extern,

+1, it should be nice to avoid unnecessary streaming of externs; as for the
latter point, wouldn't moving handling from frontends to a point in the
middle-end when the symtab is complete solve that automatically?

> and we need some sensible handling of the case where the var is only
> declared extern and omp declare target, used in some target region, but not
> actually defined anywhere in the same shared library or executable.

I think on NVPTX it yields a link error at mkoffload time.

Alexander