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?
Julian
ChangeLog
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 "" }
--
2.23.0
