https://gcc.gnu.org/bugzilla/show_bug.cgi?id=127273

--- Comment #2 from Tobias Burnus <burnus at gcc dot gnu.org> ---
OK - I messed partially up: It fails also with 'amdgcn-amdhsa',
i.e. it is likely a generic issue.

* * *

The testcase is now:

  testsuite/libgomp.c/usm_env_handling-1-self.c

(which enables ''#pragma omp requires self_maps' and includes
 testsuite/libgomp.c/usm_env_handling-1.c)

That testcase is for now XFAILED on nvptx (and will be xfailed for amdgcn as
well)

* * *

Simplified testcase [note that OpenACC rejects 'link' at compile time]

int A[1] = {1};
#pragma omp declare target link(A)
//#pragma acc declare device_resident(A)  // OK
//#pragma acc declare link(A)  // error: ‘A’ with ‘link’ clause used in
‘routine’ function

//#pragma acc routine
void
checkA (int expected, int also_ok, int other)
{
  if (expected != A[0] && other != A[0])
    __builtin_printf ("%d\n", A[0]);
  if (expected != A[0] && also_ok != A[0])
    __builtin_abort ();
}

int
main ()
{
  A[0] = 2;
  //#pragma acc parallel copyin(A)
  #pragma omp target map(always, to: A)
    checkA (2, 2, 2);
}

* * *

The failing code is the statement:

_14 = A$linkptr;

with

(gdb) p debug_bb(bb)
<bb 5> [local count: 1073741824]:
_14 = A$linkptr;
_2 = (*_14)[0];
if (_2 != expected_5(D))
  goto <bb 6>; [66.00%]
else
  goto <bb 8>; [34.00%]

* * *

The code:
  symtab_node::get (t)
is

497         gcc_checking_assert (TREE_CODE (decl) == FUNCTION_DECL
498                              || (TREE_CODE (decl) == VAR_DECL
499                                  && (TREE_STATIC (decl) || DECL_EXTERNAL 
500                                      || in_lto_p)));
501         /* Check that the mapping is sane - perhaps this check can go away,
502            but at the moment frontends tends to corrupt the mapping by
calling
503            memcpy/memset on the tree nodes.  */
504         gcc_checking_assert (!decl->decl_with_vis.symtab_node
505                              || decl->decl_with_vis.symtab_node->decl ==
decl);
506         return decl->decl_with_vis.symtab_node;


And here decl->decl_with_vis.symtab_node == NULL, which obviously will segfault
when getting dereferenced.

Reply via email to