Re: [PATCH 3/3] OpenACC dynamic data lifetimes ending within structured blocks

[Thomas Schwinge <thomas@codesourcery.com>]

Hi Julian!

On 2020-01-17T13:18:21-0800, Julian Brown <julian@codesourcery.com> wrote:
> This patch adds a new function to logically decrement the "dynamic
> reference counter" for a mapped OpenACC variable, and handles some cases
> in which that counter drops to zero inside a structured data
> block. Previously, it's likely that at least in some cases, ending a
> dynamic data lifetime in this way could behave unpredictably.
>
> Several new test cases are included.

As discussed before, all these test cases were already PASSing before any
of this thread's suggested patches (also for GCC 9), so "from a user's
point of view", all we get here are testsuite regressions:

  - 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c'
  - 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c'
  - 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c'
  - 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c'
  - 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c'
  - 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c'

(Adjusted for the version of the test cases already committed; but
already XFAILed in your original patch submission, see below.)


And: the code changes proposed here are breaking compatibility with GCC
9, such that OpenACC/Fortran code compiled with GCC 9, but running with
recent runtime libraries (common case for users, distributions) would
then terminate with: 'libgomp: cannot handle 'exit data' within data
region'.  For example, half of all 'libgomp.oacc-fortran' test cases
using OpenACC 'exit data':

  - 'libgomp.oacc-fortran/data-2.f90'
  - 'libgomp.oacc-fortran/data-3.f90'
  - 'libgomp.oacc-fortran/data-4-2.f90'
  - 'libgomp.oacc-fortran/data-4.f90'
  - 'libgomp.oacc-fortran/if-1.f90'

Even though that "code generation problem" doesn't exist with GCC 10 and
newer, we still have to maintain ABI compatibility with existing binaries
compiled compiled with GCC 9.  (Or, as a less preferred solution, arrange
so that they use host-fallback execution insted of offloading.)


Grüße
 Thomas


> This patch is strongly related to the previous two, but is somewhat of
> a separate change, and those two patches can stand alone if this one
> gets deferred.
>
> Tested alongside the previous patches in the series with offloading to NVPTX.
>
> OK?
>
> Thanks,
>
> Julian
>
> ChangeLog
>
>       libgomp/
>       * oacc-mem.c (decr_dynamic_refcount): New function.
>       (goacc_exit_datum): Call above function.
>       (goacc_exit_data_internal): Call above function.
>       * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c: New
>       test.
>       * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c:
>       Likewise.
>       * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c:
>       Likewise.
>       * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c:
>       Likewise.
>       * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c:
>       Likewise.
>       * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c:
>       Likewise.
>       * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c:
>       Likewise.
>       * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c:
>       Likewise.
> ---
>  libgomp/oacc-mem.c                            | 128 ++++++++++----
>  .../static-dynamic-lifetimes-1-lib.c          |   3 +
>  .../static-dynamic-lifetimes-1.c              | 160 ++++++++++++++++++
>  .../static-dynamic-lifetimes-6-lib.c          |   5 +
>  .../static-dynamic-lifetimes-6.c              |  46 +++++
>  .../static-dynamic-lifetimes-7-lib.c          |   5 +
>  .../static-dynamic-lifetimes-7.c              |  45 +++++
>  .../static-dynamic-lifetimes-8-lib.c          |   5 +
>  .../static-dynamic-lifetimes-8.c              |  50 ++++++
>  9 files changed, 412 insertions(+), 35 deletions(-)
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c
>
> diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
> index 783e7f363fb..f34ffa67079 100644
> --- a/libgomp/oacc-mem.c
> +++ b/libgomp/oacc-mem.c
> @@ -725,6 +725,92 @@ acc_pcopyin (void *h, size_t s)
>  #endif
>
>
> +/* Perform actions necessary to decrement the dynamic reference counter for
> +   splay tree key N.  Returns TRUE on success, or FALSE on failure (e.g. if we
> +   hit a case we can't presently handle inside a data region).  */
> +
> +static bool
> +decr_dynamic_refcount (splay_tree_key n, bool finalize)
> +{
> +  if (finalize)
> +    {
> +      if (n->refcount != REFCOUNT_INFINITY)
> +     n->refcount -= n->virtual_refcount;
> +      n->virtual_refcount = 0;
> +    }
> +
> +  if (n->virtual_refcount > 0)
> +    {
> +      if (n->refcount != REFCOUNT_INFINITY)
> +     n->refcount--;
> +      n->virtual_refcount--;
> +    }
> +  /* An initial "enter data" mapping might create a target_mem_desc (in
> +     gomp_map_vars_async via goacc_enter_datum or
> +     goacc_enter_data_internal).  In that case we have a structural
> +     reference count but a zero virtual reference count: we nevertheless
> +     want to do the "exit data" operation here.  Detect the special case
> +     using a sentinel value stored in the "prev" field, which is otherwise
> +     unused for dynamic data mappings.  */
> +  else if (n->refcount > 0
> +        && n->refcount != REFCOUNT_INFINITY
> +        && n->tgt->prev == &dyn_tgt_sentinel)
> +    {
> +      n->refcount--;
> +      /* We know n->virtual_refcount is zero here, so if we still have a
> +      non-zero n->refcount we are ending a dynamically-scoped variable
> +      lifetime in the middle of a static lifetime for the same variable.
> +      If we're not careful this results in a dangling reference.  Attempt
> +      to handle this here, if only in simple cases.  E.g.:
> +
> +        #pragma acc enter data copyin(var)
> +        #pragma acc data copy(var{, ...})
> +        {
> +          #pragma acc exit data copyout(var)
> +        }
> +
> +      Here (the "exit data"), we reattach the relevant fields of the
> +      previously dynamically-scoped target_mem_desc to the static data
> +      region's target_mem_desc, hence merging the former into the latter.
> +      The old dynamic target_mem_desc can then be freed.
> +
> +      We can't deal with static data regions that refer to existing dynamic
> +      data mappings or that introduce new static lifetimes of their own.  */
> +      if (n->refcount > 0
> +       && n->tgt->list_count == 1
> +       && n->tgt->refcount == 1)
> +     {
> +       struct goacc_thread *thr = goacc_thread ();
> +       struct target_mem_desc *tgt, *static_tgt = NULL;
> +       for (tgt = thr->mapped_data;
> +            tgt != NULL && static_tgt == NULL;
> +            tgt = tgt->prev)
> +         for (int j = 0; j < tgt->list_count; j++)
> +           if (tgt->list[j].key == n)
> +             {
> +               static_tgt = tgt;
> +               break;
> +             }
> +       if (!static_tgt
> +           || static_tgt->to_free != NULL
> +           || static_tgt->array != NULL)
> +         return false;
> +       static_tgt->to_free = n->tgt->to_free;
> +       static_tgt->array = n->tgt->array;
> +       static_tgt->tgt_start = n->tgt->tgt_start;
> +       static_tgt->tgt_end = n->tgt->tgt_end;
> +       static_tgt->to_free = n->tgt->to_free;
> +       static_tgt->refcount++;
> +       free (n->tgt);
> +       n->tgt = static_tgt;
> +     }
> +      else if (n->refcount > 0)
> +     return false;
> +    }
> +
> +  return true;
> +}
> +
>  /* Exit a dynamic mapping for a single variable.  */
>
>  static void
> @@ -767,29 +853,12 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async)
>
>    bool finalize = (kind == GOMP_MAP_DELETE
>                  || kind == GOMP_MAP_FORCE_FROM);
> -  if (finalize)
> -    {
> -      if (n->refcount != REFCOUNT_INFINITY)
> -     n->refcount -= n->virtual_refcount;
> -      n->virtual_refcount = 0;
> -    }
>
> -  if (n->virtual_refcount > 0)
> +  if (!decr_dynamic_refcount (n, finalize))
>      {
> -      if (n->refcount != REFCOUNT_INFINITY)
> -     n->refcount--;
> -      n->virtual_refcount--;
> +      gomp_mutex_unlock (&acc_dev->lock);
> +      gomp_fatal ("cannot handle delete/copyout within data region");
>      }
> -  /* An initial "enter data" mapping might create a target_mem_desc (in
> -     gomp_map_vars_async via goacc_enter_datum).  In that case we have a
> -     structural reference count but a zero virtual reference count: we
> -     nevertheless want to do the "exit data" operation here.  Detect the
> -     special case using a sentinel value stored in the "prev" field, which is
> -     otherwise unused for dynamic data mappings.  */
> -  else if (n->refcount > 0
> -        && n->refcount != REFCOUNT_INFINITY
> -        && n->tgt->prev == &dyn_tgt_sentinel)
> -    n->refcount--;
>
>    if (n->refcount == 0)
>      {
> @@ -1216,23 +1285,12 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
>           if (n == NULL)
>             continue;
>
> -         if (finalize)
> -           {
> -             if (n->refcount != REFCOUNT_INFINITY)
> -               n->refcount -= n->virtual_refcount;
> -             n->virtual_refcount = 0;
> -           }
> -
> -         if (n->virtual_refcount > 0)
> +         if (!decr_dynamic_refcount (n, finalize))
>             {
> -             if (n->refcount != REFCOUNT_INFINITY)
> -               n->refcount--;
> -             n->virtual_refcount--;
> +             /* The user is trying to do something too tricky for us.  */
> +             gomp_mutex_unlock (&acc_dev->lock);
> +             gomp_fatal ("cannot handle 'exit data' within data region");
>             }
> -         else if (n->refcount > 0
> -                  && n->refcount != REFCOUNT_INFINITY
> -                  && n->tgt->prev == &dyn_tgt_sentinel)
> -           n->refcount--;
>
>           if (copyfrom
>               && n->refcount != REFCOUNT_INFINITY
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c
> new file mode 100644
> index 00000000000..23c20d4fab7
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c
> @@ -0,0 +1,3 @@
> +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
> +/* { dg-additional-options "-DOPENACC_API" } */
> +#include "static-dynamic-lifetimes-1.c"
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c
> new file mode 100644
> index 00000000000..a743660f53e
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c
> @@ -0,0 +1,160 @@
> +/* Test transitioning of data lifetimes between static and dynamic.  */
> +
> +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
> +
> +#include <openacc.h>
> +#include <assert.h>
> +#include <stdlib.h>
> +
> +#define SIZE 1024
> +
> +void
> +f1 (void)
> +{
> +  char *block1 = (char *) malloc (SIZE);
> +
> +#ifdef OPENACC_API
> +  acc_copyin (block1, SIZE);
> +  acc_copyin (block1, SIZE);
> +#else
> +#pragma acc enter data copyin(block1[0:SIZE])
> +#pragma acc enter data copyin(block1[0:SIZE])
> +#endif
> +
> +#pragma acc data copy(block1[0:SIZE])
> +  {
> +#ifdef OPENACC_API
> +    acc_copyin (block1, SIZE);
> +#else
> +#pragma acc enter data copyin(block1[0:SIZE])
> +#endif
> +  }
> +
> +  assert (acc_is_present (block1, SIZE));
> +
> +#ifdef OPENACC_API
> +  acc_copyout (block1, SIZE);
> +  assert (acc_is_present (block1, SIZE));
> +  acc_copyout (block1, SIZE);
> +  assert (acc_is_present (block1, SIZE));
> +  acc_copyout (block1, SIZE);
> +  assert (!acc_is_present (block1, SIZE));
> +#else
> +#pragma acc exit data copyout(block1[0:SIZE])
> +  assert (acc_is_present (block1, SIZE));
> +#pragma acc exit data copyout(block1[0:SIZE])
> +  assert (acc_is_present (block1, SIZE));
> +#pragma acc exit data copyout(block1[0:SIZE])
> +  assert (!acc_is_present (block1, SIZE));
> +#endif
> +
> +  free (block1);
> +}
> +
> +void
> +f2 (void)
> +{
> +  char *block1 = (char *) malloc (SIZE);
> +
> +#ifdef OPENACC_API
> +  acc_copyin (block1, SIZE);
> +#else
> +#pragma acc enter data copyin(block1[0:SIZE])
> +#endif
> +
> +#pragma acc data copy(block1[0:SIZE])
> +  {
> +#ifdef OPENACC_API
> +    acc_copyout (block1, SIZE);
> +#else
> +#pragma acc exit data copyout(block1[0:SIZE])
> +#endif
> +    /* This should stay present until the end of the static data lifetime.  */
> +    assert (acc_is_present (block1, SIZE));
> +  }
> +
> +  assert (!acc_is_present (block1, SIZE));
> +
> +  free (block1);
> +}
> +
> +void
> +f3 (void)
> +{
> +  char *block1 = (char *) malloc (SIZE);
> +
> +#ifdef OPENACC_API
> +  acc_copyin (block1, SIZE);
> +#else
> +#pragma acc enter data copyin(block1[0:SIZE])
> +#endif
> +
> +#pragma acc data copy(block1[0:SIZE])
> +  {
> +#ifdef OPENACC_API
> +    acc_copyout (block1, SIZE);
> +    acc_copyin (block1, SIZE);
> +#else
> +#pragma acc exit data copyout(block1[0:SIZE])
> +#pragma acc enter data copyin(block1[0:SIZE])
> +#endif
> +    assert (acc_is_present (block1, SIZE));
> +  }
> +
> +  assert (acc_is_present (block1, SIZE));
> +#ifdef OPENACC_API
> +  acc_copyout (block1, SIZE);
> +#else
> +#pragma acc exit data copyout(block1[0:SIZE])
> +#endif
> +  assert (!acc_is_present (block1, SIZE));
> +
> +  free (block1);
> +}
> +
> +void
> +f4 (void)
> +{
> +  char *block1 = (char *) malloc (SIZE);
> +  char *block2 = (char *) malloc (SIZE);
> +  char *block3 = (char *) malloc (SIZE);
> +
> +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE], block3[0:SIZE])
> +  {
> +  /* The first copyin of block2 is the enclosing data region.  This
> +     "enter data" should make it live beyond the end of this region.
> +     This works, though the on-target copies of block1, block2 and block3
> +     will stay allocated until block2 is unmapped because they are bound
> +     together in a single target_mem_desc.  */
> +#ifdef OPENACC_API
> +    acc_copyin (block2, SIZE);
> +#else
> +#pragma acc enter data copyin(block2[0:SIZE])
> +#endif
> +  }
> +
> +  assert (!acc_is_present (block1, SIZE));
> +  assert (acc_is_present (block2, SIZE));
> +  assert (!acc_is_present (block3, SIZE));
> +
> +#ifdef OPENACC_API
> +  acc_copyout (block2, SIZE);
> +#else
> +#pragma acc exit data copyout(block2[0:SIZE])
> +#endif
> +  assert (!acc_is_present (block2, SIZE));
> +
> +  free (block1);
> +  free (block2);
> +  free (block3);
> +}
> +
> +int
> +main (int argc, char *argv[])
> +{
> +  f1 ();
> +  f2 ();
> +  f3 ();
> +  f4 ();
> +  return 0;
> +}
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c
> new file mode 100644
> index 00000000000..8507a0586a5
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c
> @@ -0,0 +1,5 @@
> +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
> +/* { dg-additional-options "-DOPENACC_API" } */
> +#include "static-dynamic-lifetimes-6.c"
> +/* { dg-output "libgomp: cannot handle delete/copyout within data region" } */
> +/* { dg-shouldfail "" } */
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c
> new file mode 100644
> index 00000000000..ca3b385fbcc
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c
> @@ -0,0 +1,46 @@
> +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
> +
> +#include <openacc.h>
> +#include <assert.h>
> +#include <stdlib.h>
> +
> +#define SIZE 1024
> +
> +int
> +main (int argc, char *argv[])
> +{
> +  char *block1 = (char *) malloc (SIZE);
> +  char *block2 = (char *) malloc (SIZE);
> +
> +#ifdef OPENACC_API
> +  acc_copyin (block1, SIZE);
> +  acc_copyin (block2, SIZE);
> +#else
> +#pragma acc enter data copyin(block1[0:SIZE], block2[0:SIZE])
> +#endif
> +
> +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
> +  {
> +#ifdef OPENACC_API
> +    acc_copyout (block1, SIZE);
> +    acc_copyout (block2, SIZE);
> +    /* Error output checked in static-dynamic-lifetimes-6-lib.c.  */
> +#else
> +#pragma acc exit data copyout(block1[0:SIZE], block2[0:SIZE])
> +/* We can only do this for a single dynamic data mapping at present.  */
> +/* { dg-output "libgomp: cannot handle .exit data. within data region" } */
> +/* { dg-shouldfail "" } */
> +#endif
> +    /* These should stay present until the end of the static data lifetime.  */
> +    assert (acc_is_present (block1, SIZE));
> +    assert (acc_is_present (block2, SIZE));
> +  }
> +
> +  assert (!acc_is_present (block1, SIZE));
> +  assert (!acc_is_present (block2, SIZE));
> +
> +  free (block1);
> +  free (block2);
> +
> +  return 0;
> +}
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c
> new file mode 100644
> index 00000000000..962b5926f79
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c
> @@ -0,0 +1,5 @@
> +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
> +/* { dg-additional-options "-DOPENACC_API" } */
> +#include "static-dynamic-lifetimes-7.c"
> +/* { dg-output "libgomp: cannot handle delete/copyout within data region" } */
> +/* { dg-shouldfail "" } */
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c
> new file mode 100644
> index 00000000000..dfcc7cae961
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c
> @@ -0,0 +1,45 @@
> +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
> +
> +#include <openacc.h>
> +#include <assert.h>
> +#include <stdlib.h>
> +
> +#define SIZE 1024
> +
> +int
> +main (int argc, char *argv[])
> +{
> +  char *block1 = (char *) malloc (SIZE);
> +  char *block2 = (char *) malloc (SIZE);
> +
> +#ifdef OPENACC_API
> +  acc_copyin (block1, SIZE);
> +#else
> +#pragma acc enter data copyin(block1[0:SIZE])
> +#endif
> +
> +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
> +  {
> +/* We can't attach the dynamic data mapping's (block1) target_mem_desc to the
> +   enclosing static data region here, because that region maps block2 also.  */
> +#ifdef OPENACC_API
> +    acc_copyout (block1, SIZE);
> +    /* Error output checked in static-dynamic-lifetimes-7-lib.c.  */
> +#else
> +#pragma acc exit data copyout(block1[0:SIZE])
> +/* { dg-output "libgomp: cannot handle .exit data. within data region" } */
> +/* { dg-shouldfail "" } */
> +#endif
> +    /* These should stay present until the end of the static data lifetime.  */
> +    assert (acc_is_present (block1, SIZE));
> +    assert (acc_is_present (block2, SIZE));
> +  }
> +
> +  assert (!acc_is_present (block1, SIZE));
> +  assert (!acc_is_present (block2, SIZE));
> +
> +  free (block1);
> +  free (block2);
> +
> +  return 0;
> +}
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c
> new file mode 100644
> index 00000000000..2581d7e2559
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c
> @@ -0,0 +1,5 @@
> +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
> +/* { dg-additional-options "-DOPENACC_API" } */
> +#include "static-dynamic-lifetimes-8.c"
> +/* { dg-output "libgomp: cannot handle delete/copyout within data region" } */
> +/* { dg-shouldfail "" } */
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c
> new file mode 100644
> index 00000000000..e3a64399fe9
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c
> @@ -0,0 +1,50 @@
> +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
> +
> +#include <openacc.h>
> +#include <assert.h>
> +#include <stdlib.h>
> +
> +#define SIZE 1024
> +
> +int
> +main (int argc, char *argv[])
> +{
> +  char *block1 = (char *) malloc (SIZE);
> +  char *block2 = (char *) malloc (SIZE);
> +
> +#ifdef OPENACC_API
> +  acc_copyin (block1, SIZE);
> +#else
> +#pragma acc enter data copyin(block1[0:SIZE])
> +#endif
> +
> +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
> +  {
> +#ifdef OPENACC_API
> +    acc_copyout (block1, SIZE);
> +    acc_copyin (block2, SIZE);
> +    /* Error output checked in static-dynamic-lifetimes-8-lib.c.  */
> +#else
> +#pragma acc exit data copyout(block1[0:SIZE])
> +/* { dg-output "libgomp: cannot handle .exit data. within data region" } */
> +/* { dg-shouldfail "" } */
> +#pragma acc enter data copyin(block2[0:SIZE])
> +#endif
> +    assert (acc_is_present (block1, SIZE));
> +    assert (acc_is_present (block2, SIZE));
> +  }
> +
> +  assert (!acc_is_present (block1, SIZE));
> +  assert (acc_is_present (block2, SIZE));
> +#ifdef OPENACC_API
> +  acc_copyout (block2, SIZE);
> +#else
> +#pragma acc exit data copyout(block2[0:SIZE])
> +#endif
> +  assert (!acc_is_present (block2, SIZE));
> +
> +  free (block1);
> +  free (block2);
> +
> +  return 0;
> +}
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter