| Message ID | 4673a5070087e465f6dd123715d409b35b875ca1.1579292772.git.julian@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | Mixed static/dynamic data lifetimes with OpenACC (PR92843) | expand |
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
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; +}