Hi,

if we compile the testcase pr84592-2.c from the patch:
...
#include <stdlib.h>

int
main (void)
{
   int n[1];

   n[0] = 3;

#pragma omp target
   {
     static int test[4] = { 1, 2, 3, 4 };
     n[0] += test[n[0]];
   }

   if (n[0] != 7)
     abort ();

   return 0;
}
...

for nvptx offloading, we run into an assert:
...
lto1: internal compiler error: in input_varpool_node, at lto-cgraph.c:1424
0x959ebb input_varpool_node
         gcc/lto-cgraph.c:1422
0x959ebb input_cgraph_1
         gcc/lto-cgraph.c:1544
0x959ebb input_symtab()
         gcc/lto-cgraph.c:1858
0x5aceac read_cgraph_and_symbols
         gcc/lto/lto.c:2891
0x5aceac lto_main()
         gcc/lto/lto.c:3356
...

The assert we run into is:
...
1422      gcc_assert (flag_ltrans
1423                  || (!node->in_other_partition
1424                      && !node->used_from_other_partition));
...

where node is:
...
(gdb) call debug_generic_expr (node.decl)
test
...

and the reason the assert triggers is:
...
(gdb) p node.in_other_partition
$1 = 1
...

AFAIU, what this means is that the variable test is placed in a 
different partition than the offloading function main._omp_fn.0 that 
uses the variable.


I looked at where global variables are put into offload_vars, and found 
that that happens in varpool_node::get_create:
...
   if ((flag_openacc || flag_openmp)
       && lookup_attribute ("omp declare target",
                            DECL_ATTRIBUTES (decl)))
     {
       node->offloadable = 1;
       if (ENABLE_OFFLOADING && !DECL_EXTERNAL (decl))
         {
           g->have_offload = true;
           if (!in_lto_p)
             vec_safe_push (offload_vars, decl);
         }
     }

...

The patch fixes the ICE there by marking the varpool_node test as 
offloadable as well.

Build and reg-tested libgomp on x86_64 with nvptx accelerator.
Bootstrapped and reg-tested on x86_64.

OK for stage4 trunk?

Thanks,
- Tom