Re: [PATCH 1/2] [OpenACC] Refuse update/copyout for blocks with attached pointers

[Thomas Schwinge <thomas@codesourcery.com>]

Hi Julian!

Quoting your parent email:

On 2020-06-22T05:14:42-0700, Julian Brown <julian@codesourcery.com> wrote:
> Investigating PR95590, I realised that we can do better at diagnosing
> some potentially troublesome usage of OpenACC "attach" behaviour, namely
> updating blocks with attached pointers. Updating either the host copy
> or device copy of such a block is problematic -- for a host update,
> the host may get a clobbered (device) version of a host pointer in its
> local version of the block (e.g. struct). A device update may clobber
> an attached device pointer with a host pointer.

ACK.

> The spec text (OpenACC 3.0, "2.6.8. Attachment Counter") covering this
> case is:
>
>   "Pointer members of structs, classes, or derived types in device
>    or host memory can be overwritten due to update directives or API
>    routines. It is the user’s responsibility to ensure that the pointers
>    have the appropriate values before or after the data movement in
>    either direction. The behavior of the program is undefined if any
>    of the pointer members are attached when an update of a composite
>    variable is performed."
>
> The first patch in this series addresses that paragraph by making
> such updates (as well as copyouts, similarly) be runtime errors.

Hmm.  But why do you say "addresses [...] by making [...] be runtime
errors" if the specification text *explicitly* states ("It is the user's
responsibility") that doing such things invokes undefined behavior, and
thus a user must not do that.  (Here, the undefined behavior is: copying
of host vs. device pointers -- I wouldn't assume (user), respectively
imply (implementor) anything worse?)

It's of couse good if we can (without much overhead) be helpful to the
user (your proposed runtime error), but I want to make sure that I'm
correctly understanding your rationale here.

On 2020-06-22T05:14:43-0700, Julian Brown <julian@codesourcery.com> wrote:
> As mentioned in the parent email, this patch adds diagnostics for
> probably-broken code that updates (host/device) or copies-out blocks
> that still have attached pointers.  Several new tests have been added.
>
> OK?

I so far haven't managed to really convince myself that we want to incur
this overhead here.  (I suppose it's not too much overhead.)  I may
re-consider this still.

I suppose we can put this onto the backburner -- nothing else
functionally depends on this?


Assuming this checking does get installed (and enabled by default), I had
the idea that we may (rather easily?) add a flag variable (ICV;
initialized from an environment variable) to guard this checking
behavior?  I suppose we may now have a few libgomp testcases that
actually do use 'acc_update_self' etc. to read out pointer values from
visible device copies, and verify these, which wouldn't work any longer
with that checking enabled.  Such tests could then 'dg-set-target-env-var
"GOMP_ATTACH_CHECKING" "0"' (better name is desirable), and have one
variant with and one variant without the checking.


Grüße
 Thomas


>       libgomp/
>       * oacc-mem.c (update_dev_host): Raise error on update of block with
>       attached pointers.
>       (goacc_exit_data_internal): Raise error on copyout of block with
>       attached pointers.
>       * target.c (gomp_unmap_vars_internal): Likewise.
>       * testsuite/libgomp.oacc-c-c++-common/copyback-attached-dynamic-1.c:
>       New test.
>       * testsuite/libgomp.oacc-c-c++-common/delete-attached-dynamic-1.c:
>       New test.
>       * testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-1.c:
>       New test.
>       * testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-2.c:
>       New test.
>       * testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-3.c:
>       New test.
>       * testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-1.c:
>       New test.
>       * testsuite/libgomp.oacc-c-c++-common/update-attached.c: New test.
>       * testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90: Update
>       for new diagnostic.
> ---
>  libgomp/oacc-mem.c                            | 42 ++++++++++++++++---
>  libgomp/target.c                              | 27 +++++++++---
>  .../copyback-attached-dynamic-1.c             | 31 ++++++++++++++
>  .../copyback-attached-structural-1.c          | 30 +++++++++++++
>  .../copyback-attached-structural-2.c          | 31 ++++++++++++++
>  .../copyback-attached-structural-3.c          | 26 ++++++++++++
>  .../delete-attached-dynamic-1.c               | 26 ++++++++++++
>  .../delete-attached-structural-1.c            | 25 +++++++++++
>  .../delete-attached-structural-2.c            | 26 ++++++++++++
>  .../update-attached-1.c                       | 33 +++++++++++++++
>  .../deep-copy-6-no_finalize.F90               |  6 +--
>  11 files changed, 290 insertions(+), 13 deletions(-)
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-dynamic-1.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-1.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-2.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-3.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-dynamic-1.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-1.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-2.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/update-attached-1.c
>
> diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
> index 1816b06bf2d..cf054f14b12 100644
> --- a/libgomp/oacc-mem.c
> +++ b/libgomp/oacc-mem.c
> @@ -865,6 +865,23 @@ update_dev_host (int is_dev, void *h, size_t s, int async)
>        gomp_fatal ("[%p,%d] is not mapped", h, (int)s);
>      }
>
> +  if (n->aux && n->aux->attach_count)
> +    {
> +      size_t nptrs = (n->host_end - n->host_start + sizeof (void *) - 1)
> +                  / sizeof (void *);
> +      for (size_t i = 0; i < nptrs; i++)
> +     if (n->aux->attach_count[i] > 0)
> +       {
> +         gomp_mutex_unlock (&acc_dev->lock);
> +         if (is_dev)
> +           gomp_fatal ("[%p,+%d] device update would overwrite attached "
> +                       "pointers", h, (int) s);
> +         else
> +           gomp_fatal ("host update from block [%p,+%d] with attached "
> +                       "pointers", h, (int) s);
> +       }
> +    }
> +
>    d = (void *) (n->tgt->tgt_start + n->tgt_offset
>               + (uintptr_t) h - n->host_start);
>
> @@ -1329,11 +1346,26 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
>           if (copyfrom
>               && n->refcount != REFCOUNT_INFINITY
>               && (kind != GOMP_MAP_FROM || n->refcount == 0))
> -           gomp_copy_dev2host (acc_dev, aq, (void *) cur_node.host_start,
> -                               (void *) (n->tgt->tgt_start + n->tgt_offset
> -                                         + cur_node.host_start
> -                                         - n->host_start),
> -                               cur_node.host_end - cur_node.host_start);
> +           {
> +             if (n->aux && n->aux->attach_count)
> +               {
> +                 size_t nptrs = (n->host_end - n->host_start
> +                                 + sizeof (void *) - 1) / sizeof (void *);
> +                 for (size_t j = 0; j < nptrs; j++)
> +                   if (n->aux->attach_count[j] > 0)
> +                     {
> +                       gomp_mutex_unlock (&acc_dev->lock);
> +                       gomp_fatal ("copyout of block [%p,+%d] with "
> +                                   "attached pointers", hostaddrs[i],
> +                                   (int) size);
> +                     }
> +               }
> +             gomp_copy_dev2host (acc_dev, aq, (void *) cur_node.host_start,
> +                                 (void *) (n->tgt->tgt_start + n->tgt_offset
> +                                           + cur_node.host_start
> +                                           - n->host_start),
> +                                 cur_node.host_end - cur_node.host_start);
> +           }
>
>           if (n->refcount == 0)
>             {
> diff --git a/libgomp/target.c b/libgomp/target.c
> index badc254a777..db6f56a8ff8 100644
> --- a/libgomp/target.c
> +++ b/libgomp/target.c
> @@ -1649,11 +1649,28 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
>
>        if ((do_unmap && do_copyfrom && tgt->list[i].copy_from)
>         || tgt->list[i].always_copy_from)
> -     gomp_copy_dev2host (devicep, aq,
> -                         (void *) (k->host_start + tgt->list[i].offset),
> -                         (void *) (k->tgt->tgt_start + k->tgt_offset
> -                                   + tgt->list[i].offset),
> -                         tgt->list[i].length);
> +     {
> +       if (k->aux && k->aux->attach_count)
> +         {
> +           size_t nptrs = (k->host_end - k->host_start
> +                           + sizeof (void *) - 1) / sizeof (void *);
> +           for (size_t j = 0; j < nptrs; j++)
> +             if (k->aux->attach_count[j] > 0)
> +               {
> +                 gomp_mutex_unlock (&devicep->lock);
> +                 gomp_fatal ("copyout of block [%p,+%d] with "
> +                             "attached pointers",
> +                             (void *) (k->host_start + tgt->list[i].offset),
> +                             (int) (k->host_end - k->host_start));
> +               }
> +         }
> +       gomp_copy_dev2host (devicep, aq,
> +                           (void *) (k->host_start + tgt->list[i].offset),
> +                           (void *) (k->tgt->tgt_start + k->tgt_offset
> +                                     + tgt->list[i].offset),
> +                           tgt->list[i].length);
> +     }
> +
>        if (do_unmap)
>       {
>         struct target_mem_desc *k_tgt = k->tgt;
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-dynamic-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-dynamic-1.c
> new file mode 100644
> index 00000000000..bc4e297fa6f
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-dynamic-1.c
> @@ -0,0 +1,31 @@
> +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
> +
> +#include <stdlib.h>
> +
> +struct mystruct {
> +  int *arr;
> +};
> +
> +int
> +main (int argc, char *argv[])
> +{
> +  int localarray[1024];
> +  struct mystruct s;
> +  s.arr = localarray;
> +
> +  #pragma acc enter data copyin(s)
> +
> +  #pragma acc data copy(s.arr[0:1024])
> +  {
> +    /* This directive does one too many attachments: it should fail when we try
> +       to do the copyout below.  */
> +    #pragma acc enter data attach(s.arr)
> +    /* { dg-output "copyout of block \\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] with attached pointers" } */
> +  }
> +
> +  #pragma acc exit data copyout(s)
> +
> +  return 0;
> +}
> +
> +/* { dg-shouldfail "" } */
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-1.c
> new file mode 100644
> index 00000000000..7846c8c717c
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-1.c
> @@ -0,0 +1,30 @@
> +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
> +
> +#include <stdlib.h>
> +
> +struct mystruct {
> +  int *arr;
> +};
> +
> +int
> +main (int argc, char *argv[])
> +{
> +  int localarray[1024];
> +  struct mystruct s;
> +  s.arr = localarray;
> +
> +  #pragma acc data copy(s)
> +  {
> +    #pragma acc data copy(s.arr[0:1024])
> +    {
> +      /* This directive does one too many attachments: it should fail when we try
> +      to do the copyout below.  */
> +      #pragma acc enter data attach(s.arr)
> +      /* { dg-output "copyout of block \\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] with attached pointers" } */
> +    }
> +  }
> +
> +  return 0;
> +}
> +
> +/* { dg-shouldfail "" } */
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-2.c
> new file mode 100644
> index 00000000000..bffa06eb725
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-2.c
> @@ -0,0 +1,31 @@
> +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
> +
> +#include <stdlib.h>
> +
> +struct mystruct {
> +  int *arr;
> +};
> +
> +int
> +main (int argc, char *argv[])
> +{
> +  int localarray[1024];
> +  struct mystruct s;
> +  s.arr = localarray;
> +
> +  #pragma acc enter data copyin(localarray[0:1024])
> +
> +  #pragma acc data copy(s)
> +  {
> +    /* This directive does one too many attachments: it should fail when we try
> +       to do the copyout below.  */
> +    #pragma acc enter data attach(s.arr)
> +    /* { dg-output "copyout of block \\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] with attached pointers" } */
> +  }
> +
> +  #pragma acc exit data delete(localarray[0:1024])
> +
> +  return 0;
> +}
> +
> +/* { dg-shouldfail "" } */
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-3.c
> new file mode 100644
> index 00000000000..4b21677af09
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-3.c
> @@ -0,0 +1,26 @@
> +#include <stdlib.h>
> +
> +struct mystruct {
> +  int *arr;
> +};
> +
> +int
> +main (int argc, char *argv[])
> +{
> +  int localarray[1024];
> +  struct mystruct s;
> +  s.arr = localarray;
> +
> +  #pragma acc enter data copyin(localarray[0:1024])
> +
> +  #pragma acc data copy(s)
> +  {
> +    /* Here the attach and detach balance: this should work.  */
> +    #pragma acc enter data attach(s.arr)
> +    #pragma acc exit data detach(s.arr)
> +  }
> +
> +  #pragma acc exit data delete(localarray[0:1024])
> +
> +  return 0;
> +}
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-dynamic-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-dynamic-1.c
> new file mode 100644
> index 00000000000..e074d507fb2
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-dynamic-1.c
> @@ -0,0 +1,26 @@
> +#include <stdlib.h>
> +
> +struct mystruct {
> +  int *arr;
> +};
> +
> +int
> +main (int argc, char *argv[])
> +{
> +  int localarray[1024];
> +  struct mystruct s;
> +  s.arr = localarray;
> +
> +  #pragma acc enter data copyin(s)
> +
> +  #pragma acc data copy(s.arr[0:1024])
> +  {
> +    /* We delete 's' from the target below: this extra attachment is not
> +       dangerous and we do not raise an error.  */
> +    #pragma acc enter data attach(s.arr)
> +  }
> +
> +  #pragma acc exit data delete(s)
> +
> +  return 0;
> +}
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-1.c
> new file mode 100644
> index 00000000000..e675762ecd8
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-1.c
> @@ -0,0 +1,25 @@
> +#include <stdlib.h>
> +
> +struct mystruct {
> +  int *arr;
> +};
> +
> +int
> +main (int argc, char *argv[])
> +{
> +  int localarray[1024];
> +  struct mystruct s;
> +  s.arr = localarray;
> +
> +  #pragma acc data copyin(s)
> +  {
> +    #pragma acc data copy(s.arr[0:1024])
> +    {
> +      /* This directive does one too many attachments: it should fail when we try
> +      to do the copyout below.  */
> +      #pragma acc enter data attach(s.arr)
> +    }
> +  }
> +
> +  return 0;
> +}
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-2.c
> new file mode 100644
> index 00000000000..d2095255ad3
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-2.c
> @@ -0,0 +1,26 @@
> +#include <stdlib.h>
> +
> +struct mystruct {
> +  int *arr;
> +};
> +
> +int
> +main (int argc, char *argv[])
> +{
> +  int localarray[1024];
> +  struct mystruct s;
> +  s.arr = localarray;
> +
> +  #pragma acc enter data copyin(localarray[0:1024])
> +
> +  #pragma acc data copyin(s)
> +  {
> +    /* We only try to copy in: the extra attachment we're left over with is not
> +       harmful and we don't raise an error.  */
> +    #pragma acc enter data attach(s.arr)
> +  }
> +
> +  #pragma acc exit data delete(localarray[0:1024])
> +
> +  return 0;
> +}
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/update-attached-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/update-attached-1.c
> new file mode 100644
> index 00000000000..9f60bfa56f4
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/update-attached-1.c
> @@ -0,0 +1,33 @@
> +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
> +
> +#include <stdlib.h>
> +
> +struct mystruct {
> +  int *arr;
> +};
> +
> +int
> +main (int argc, char *argv[])
> +{
> +  int localarray[1024];
> +  int localarray2[1024];
> +  struct mystruct s;
> +  s.arr = localarray;
> +
> +  #pragma acc enter data copyin(s)
> +
> +  #pragma acc data copy(s.arr[0:1024])
> +  {
> +    s.arr = localarray2;
> +    /* This update is dangerous because we have attached pointers: raise an
> +       error.  */
> +    #pragma acc update device(s)
> +    /* { dg-output "\\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] device update would overwrite attached pointers" } */
> +  }
> +
> +  #pragma acc exit data delete(s)
> +
> +  return 0;
> +}
> +
> +/* { dg-shouldfail "" } */
> diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90
> index ad8da71d7c9..355a381b625 100644
> --- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90
> @@ -8,7 +8,7 @@
>  ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
>  ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
>
> -! Without the finalize, we do not detach properly so the host sees a device
> -! pointer, and fails with this STOP code.
> -! { dg-output "STOP 7(\n|\r\n|\r)+" }
> +! Without the finalize, we do not detach properly and raise an error on attempting
> +! the copyout.
> +! { dg-output ".*copyout of block \\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] with attached pointers(\n|\r\n|\r)+" }
>  ! { dg-shouldfail "" }
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter