[Bug middle-end/95270] OpenACC 'enter data attach' looks up target memory object displaced by pointer size

[tschwinge at gcc dot gnu.org]

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

Thomas Schwinge <tschwinge at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
           Assignee|unassigned at gcc dot gnu.org      |jules at gcc dot gnu.org
             Status|UNCONFIRMED                 |ASSIGNED
   Target Milestone|---                         |10.0
            Summary|OpenACC 'enter data         |OpenACC 'enter data attach'
                   |attach(data_p)' fails for   |looks up target memory
                   |'int *data_p'               |object displaced by pointer
                   |                            |size
   Last reconfirmed|                            |2020-06-08
     Ever confirmed|0                           |1

--- Comment #1 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
I just hit this one again, started to debug it once again (until I finally
remembered this existing issue...), and cried a little bit.  ;'-|

Julian, this is a bug, right?  During gimplification, or in the front end(s)? 
Please have a look.

This issue is also visible in the (current) 'omplower' tree dump scans in
'c-c++-common/goacc/mdc-1.c': expecting 'bias: 8'.  (Pointer size, by the way;
cross-checked with '-m32' where we get 'bias: 4'.)

I'd thought I remember that I may work around this by wrapping pointers as
'struct' members, which would then 'attach' correctly, without size/bias
confusion, but apparently that's not true; the attached 'mdc-2b_.c' fails in
the same way even if altered as follows:

    -  int *data_p = &data;
    +#if 0
    +  int *data_p;
    +#else
    +  struct s {
    +    int *data_p;
    +  } s;
    +# define data_p s.data_p
    +#endif
    +  data_p = &data;

What you instead have to do is to make sure that your target memory object
(here: 'data') is bigger than the bias (pointer size):

    -  int data;
    +  struct { char d[sizeof (void *) + 1]; } data;

Then it works, in all variants I've now posted.

Execution test cases that would ("might") exercise/notice this behavior:

  - 'libgomp.oacc-c-c++-common/deep-copy-2.c' -- works despite the bug, because
the '#pragma acc data copy(s.a[0:10])' already synthesizes an implicit
'map(attach:s.a [bias: 0])' (correct) and thus the incorrect 'map(attach:s.a
[bias: 8])' synthesized for the inner '#pragma acc parallel loop attach(s.a)'
doesn't cause a problem, as 'bias' is only used in the first call of the
respective 'gomp_attach_pointer'.
  - 'libgomp.oacc-c-c++-common/deep-copy-4.c' -- works despite the bug, because
for a struct bigger than the bias ('struct node' here), a 'splay_tree_lookup'
with start offset by bias will then still look up the expected object,
apparently.  With the 'val'/'sum' business removed, this test case quickly
fails: 'libgomp: pointer target not mapped for attach'.
  - 'libgomp.oacc-fortran/deep-copy-6.f90' -- correctly synthesizes
'map(attach:var.a [bias: 0])' for '!$acc enter data attach(var%a)'.  However,
for the earlier '!$acc enter data copyin(var%a(5:n - 5), var%b(5:n - 5))'
synthesizes the two 'gomp_attach_pointer' calls both with 'bias = 16' -- seems
to work, but is that correct?
  - 'libgomp.oacc-fortran/deep-copy-8.f90' -- correctly synthesizes
'map(attach:var.a [bias: 0])' for '!$acc enter data attach(var%a)'.

That's all with an explicit 'attach' clause, folks.