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.
