[PATCH] Fix ICE for static vars in offloaded functions

Tom de Vries Tom_deVries@mentor.com
Wed Mar 7 13:20:00 GMT 2018


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

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
0x959ebb input_cgraph_1
0x959ebb input_symtab()
0x5aceac read_cgraph_and_symbols
0x5aceac lto_main()

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)

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;
           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?

- Tom
-------------- next part --------------
A non-text attachment was scrubbed...
Name: 0001-Fix-ICE-for-static-vars-in-offloaded-functions.patch
Type: text/x-patch
Size: 3559 bytes
Desc: not available
URL: <http://gcc.gnu.org/pipermail/gcc-patches/attachments/20180307/1ac2467c/attachment.bin>

More information about the Gcc-patches mailing list