@@ -975,7 +975,8 @@ struct target_mem_desc {
uintptr_t tgt_end;
/* Handle to free. */
void *to_free;
- /* Previous target_mem_desc. */
+ /* Previous target_mem_desc. Also used in OpenACC to indicate that this
+ target_mem_desc is used only for an "enter data" mapping. */
struct target_mem_desc *prev;
/* Number of items in following list. */
size_t list_count;
@@ -165,6 +165,8 @@ bool _goacc_profiling_setup_p (struct goacc_thread *,
void goacc_profiling_dispatch (acc_prof_info *, acc_event_info *,
acc_api_info *);
+extern void goacc_mark_dynamic (struct target_mem_desc *);
+
#ifdef HAVE_ATTRIBUTE_VISIBILITY
# pragma GCC visibility pop
#endif
@@ -550,6 +550,24 @@ acc_unmap_data (void *h)
}
}
+/* Indicate (via storing its address in the "prev" field) a target_mem_desc
+ that is used for an "enter data" mapping. */
+const static struct target_mem_desc dyn_tgt_sentinel;
+
+/* Mark TGT as the "initial" target_mem_desc created by a dynamic data mapping
+ (acc_create, acc_copyin or an "enter data" directive). For such mappings,
+ to start with, we have a splay tree key with a reference count of 1 and a
+ virtual reference count of 0 (linking to this target_mem_desc). Without
+ this marking, such a mapping is indistinguishable from a target_mem_desc
+ created by e.g. a lexically-scoped "acc data" region, but the difference is
+ important if acc_copyout, acc_delete (etc.) or an "exit data" directive is
+ used to end the data lifetime. */
+
+void
+goacc_mark_dynamic (struct target_mem_desc *tgt)
+{
+ tgt->prev = (struct target_mem_desc *) &dyn_tgt_sentinel;
+}
/* Enter dynamic mapping for a single datum. Return the device pointer. */
@@ -613,8 +631,14 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async)
goacc_aq aq = get_goacc_asyncqueue (async);
- gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds,
- true, GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+ struct target_mem_desc *tgt;
+ tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes,
+ kinds, true, GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+
+ /* Mark non-NULL target_mem_descs returned here specially: see comment in
+ goacc_exit_datum. */
+ if (tgt)
+ goacc_mark_dynamic (tgt);
gomp_mutex_lock (&acc_dev->lock);
n = lookup_host (acc_dev, hostaddrs[0], sizes[0]);
@@ -756,7 +780,15 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async)
n->refcount--;
n->virtual_refcount--;
}
- else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY)
+ /* 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)
@@ -1081,11 +1113,12 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
dump_mappings ((group_last - i) + 1, &hostaddrs[i], &sizes[i], &kinds[i]);
#endif
- gomp_map_vars_async (acc_dev, aq,
- (group_last - i) + 1,
- &hostaddrs[i], NULL,
- &sizes[i], &kinds[i], true,
- GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+ struct target_mem_desc *tgt;
+ tgt = gomp_map_vars_async (acc_dev, aq, (group_last - i) + 1,
+ &hostaddrs[i], NULL, &sizes[i], &kinds[i],
+ true, GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+ if (tgt)
+ goacc_mark_dynamic (tgt);
i = group_last;
}
@@ -1196,7 +1229,9 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
n->refcount--;
n->virtual_refcount--;
}
- else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY)
+ else if (n->refcount > 0
+ && n->refcount != REFCOUNT_INFINITY
+ && n->tgt->prev == &dyn_tgt_sentinel)
n->refcount--;
if (copyfrom
@@ -1629,6 +1629,8 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
k->refcount == 1, NULL);
}
+ bool have_virtual_refs = false, all_refs_virtual = true;
+
for (i = 0; i < tgt->list_count; i++)
{
splay_tree_key k = tgt->list[i].key;
@@ -1636,21 +1638,21 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
continue;
bool do_unmap = false;
- if (k->tgt == tgt
- && k->virtual_refcount > 0
- && k->refcount != REFCOUNT_INFINITY)
- {
- k->virtual_refcount--;
- k->refcount--;
- }
- else if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
+ if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
k->refcount--;
else if (k->refcount == 1)
{
k->refcount--;
- do_unmap = true;
+ if (k->virtual_refcount == 0)
+ do_unmap = true;
}
+ if (k->virtual_refcount > 0 && k->refcount == k->virtual_refcount)
+ have_virtual_refs = true;
+
+ if (k->refcount != k->virtual_refcount)
+ all_refs_virtual = false;
+
if ((do_unmap && do_copyfrom && tgt->list[i].copy_from)
|| tgt->list[i].always_copy_from)
gomp_copy_dev2host (devicep, aq,
@@ -1670,6 +1672,42 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
}
}
+ if (have_virtual_refs)
+ {
+ /* If we have a construct such as this:
+
+ #pragma acc data copy(var1)
+ {
+ #pragma acc enter data copyin(var1)
+ }
+
+ The dynamic data lifetime entered in the middle of the static
+ data lifetime extends beyond the static lifetime. Adjust
+ references and the target descriptor here (the end of the static
+ region) to make it seem like we did "enter data" on the data to
+ start with.
+
+ We can't do this adjustment if the data construct refers to other
+ variables too. */
+ if (!all_refs_virtual)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("cannot handle create/copyin/'enter data' within data "
+ "region");
+ }
+
+ for (i = 0; i < tgt->list_count; i++)
+ {
+ splay_tree_key k = tgt->list[i].key;
+ if (k == NULL || k->virtual_refcount == 0)
+ continue;
+
+ if (k->refcount == k->virtual_refcount)
+ k->virtual_refcount--;
+ }
+ goacc_mark_dynamic (tgt);
+ }
+
if (aq)
devicep->openacc.async.queue_callback_func (aq, gomp_unref_tgt_void,
(void *) tgt);
new file mode 100644
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-2.c"
new file mode 100644
@@ -0,0 +1,166 @@
+/* Test nested dynamic/static data mappings. */
+
+/* { 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);
+
+#pragma acc data copy(block1[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+
+ 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
+
+ assert (!acc_is_present (block1, SIZE));
+
+ free (block1);
+}
+
+void
+f3 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#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);
+
+#pragma acc data copy(block1[0: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_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+#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
+f5 (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_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
+ }
+#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);
+}
+
+int
+main (int argc, char *argv[])
+{
+ f1 ();
+ f2 ();
+ f3 ();
+ f4 ();
+ f5 ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-3.c"
new file mode 100644
@@ -0,0 +1,183 @@
+/* Test nested dynamic/static data mappings (multiple blocks on data
+ regions). */
+
+/* { 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);
+ char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f2 (void)
+{
+ 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);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f3 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+ char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyin (block2, SIZE);
+ acc_copyout (block2, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc enter data copyin(block2[0:SIZE])
+#pragma acc exit data copyout(block2[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f4 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+ char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0: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_copyin (block2, SIZE);
+ acc_copyout (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+ }
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f5 (void)
+{
+ 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_copyin (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+ }
+#ifdef OPENACC_API
+ acc_copyout (block2, SIZE);
+#else
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+ }
+
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+int
+main (int argc, char *argv[])
+{
+ f1 ();
+ f2 ();
+ f3 ();
+ f4 ();
+ f5 ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,6 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-4.c"
+
+/* { dg-output "libgomp: cannot handle create/copyin/.enter data. within data region" } */
+/* { dg-shouldfail "" } */
new file mode 100644
@@ -0,0 +1,71 @@
+/* { 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);
+ char *block3 = (char *) malloc (SIZE);
+
+ /* Doing this twice ensures that we have a non-zero virtual refcount. Make
+ sure that works too. */
+#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], 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. */
+#ifdef OPENACC_API
+ acc_copyin (block2, SIZE);
+ /* Error output checked in static-dynamic-lifetimes-4-lib.c. */
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+ /* ...except that doesn't work at present because it would mean the dynamic
+ data region would get entangled with the static data region's
+ target_mem_desc that has mappings for each of block1, block2 and block3.
+ Check for runtime error. */
+ /* { dg-output "libgomp: cannot handle create/copyin/.enter data. within data region" } */
+ /* { dg-shouldfail "" } */
+#endif
+ }
+
+ assert (acc_is_present (block1, SIZE));
+ assert (acc_is_present (block2, SIZE));
+ assert (!acc_is_present (block3, 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 (block2, SIZE);
+ assert (!acc_is_present (block2, 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(block2[0:SIZE])
+ assert (!acc_is_present (block2, SIZE));
+#endif
+
+ free (block1);
+ free (block2);
+ free (block3);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,6 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-5.c"
+
+/* { dg-output "libgomp: cannot handle create/copyin/.enter data. within data region" } */
+/* { dg-shouldfail "" } */
new file mode 100644
@@ -0,0 +1,63 @@
+/* { 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);
+ char *block3 = (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], 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. */
+#ifdef OPENACC_API
+ acc_copyin (block2, SIZE);
+ /* Error output checked in static-dynamic-lifetimes-5-lib.c. */
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+ /* ...except that doesn't work at present because it would mean the dynamic
+ data region would get entangled with the static data region's
+ target_mem_desc that has mappings for each of block1, block2 and block3.
+ Check for runtime error. */
+ /* { dg-output "libgomp: cannot handle create/copyin/.enter data. within data region" } */
+ /* { dg-shouldfail "" } */
+#endif
+ }
+
+ assert (acc_is_present (block1, SIZE));
+ assert (acc_is_present (block2, SIZE));
+ assert (!acc_is_present (block3, SIZE));
+
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+ assert (!acc_is_present (block1, SIZE));
+
+ acc_copyout (block2, SIZE);
+ assert (!acc_is_present (block2, SIZE));
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+ assert (!acc_is_present (block1, SIZE));
+
+#pragma acc exit data copyout(block2[0:SIZE])
+ assert (!acc_is_present (block2, SIZE));
+#endif
+
+ free (block1);
+ free (block2);
+ free (block3);
+
+ return 0;
+}