Hi,
This patch fixes an issue I noticed when investigating an answer
for Thomas's question about device pointer return values in:
https://gcc.gnu.org/ml/gcc-patches/2019-10/msg02260.html
It looks to me like the return value for the present case is wrong in
the existing code: in case of a acc_pcopyin or similar call that refers
to a subarray of a larger block already mapped on the target, the
device pointer return value will be the start of the larger block, not
of the subarray being copied.
The attached patch corrects this issue, and also relaxes a restriction
on acc_delete, acc_copyout (etc.) to allow them to unmap/copyout
subarrays of a larger block already present on the target. There's no
particular reason to disallow that, as far as I can tell. This is
necessary to allow the new tests included with this patch to pass, and
a couple of existing "shouldfail" tests no longer fail, and have been
adjusted accordingly. It's still an error to try to copy data beyond
the bounds of a mapped block, and other existing tests cover those
cases.
The calculation for the return value for the non-present case of
present_create_copy has also been adjusted in anticipation of a new
version of the above-linked patch.
Tested with offloading to nvptx. OK for trunk?
Julian
ChangeLog
libgomp/
* oacc-mem.c (present_create_copy): Fix device pointer return value in
case of "present" subarray. Use tgt->tgt_start instead of tgt->to_free
in non-present/create case.
(delete_copyout): Change error condition to detect only copies outside
of mapped block. Adjust error message accordingly.
* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/lib-20.c: Adjust expected error
message.
* testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-22.c: Allow test to pass now.
* testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise.
commit 00607b06c8e506b0f0744a230856e1e8776633c3
Author: Julian Brown <jul...@codesourcery.com>
Date: Thu Nov 7 14:24:49 2019 -0800
OpenACC "present" subarrays: runtime API return value and unmapping fixes
libgomp/
* oacc-mem.c (present_create_copy): Fix device pointer return value in
case of "present" subarray. Use tgt->tgt_start instead of tgt->to_free
in non-present/create case.
(delete_copyout): Change error condition to fail only on copies outside
of mapped block. Adjust error message accordingly.
* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/lib-20.c: Adjust expected error
message.
* testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-22.c: Allow test to pass now.
* testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise.
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 2f271009fb8..0a41f11210c 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -535,7 +535,7 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
if (n)
{
/* Present. */
- d = (void *) (n->tgt->tgt_start + n->tgt_offset);
+ d = (void *) (n->tgt->tgt_start + n->tgt_offset + h - n->host_start);
if (!(f & FLAG_PRESENT))
{
@@ -584,7 +584,7 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
gomp_mutex_lock (&acc_dev->lock);
- d = tgt->to_free;
+ d = (void *) tgt->tgt_start;
tgt->prev = acc_dev->openacc.data_environ;
acc_dev->openacc.data_environ = tgt;
@@ -669,7 +669,6 @@ acc_pcopyin (void *h, size_t s)
static void
delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
{
- size_t host_size;
splay_tree_key n;
void *d;
struct goacc_thread *thr = goacc_thread ();
@@ -703,13 +702,12 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
d = (void *) (n->tgt->tgt_start + n->tgt_offset
+ (uintptr_t) h - n->host_start);
- host_size = n->host_end - n->host_start;
-
- if (n->host_start != (uintptr_t) h || host_size != s)
+ if ((uintptr_t) h < n->host_start || (uintptr_t) h + s > n->host_end)
{
+ size_t host_size = n->host_end - n->host_start;
gomp_mutex_unlock (&acc_dev->lock);
- gomp_fatal ("[%p,%d] surrounds2 [%p,+%d]",
- (void *) n->host_start, (int) host_size, (void *) h, (int) s);
+ gomp_fatal ("[%p,+%d] outside mapped block [%p,+%d]",
+ (void *) h, (int) s, (void *) n->host_start, (int) host_size);
}
if (n->refcount == REFCOUNT_INFINITY)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c
new file mode 100644
index 00000000000..bee0b10ca7b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c
@@ -0,0 +1,28 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <stdlib.h>
+#include <assert.h>
+#include <stdint.h>
+
+int main (int argc, char* argv[])
+{
+ char* myblock = malloc (1024);
+ int i;
+ void *dst;
+ for (i = 0; i < 1024; i++)
+ myblock[i] = i;
+ dst = acc_copyin (myblock, 1024);
+ for (i = 0; i < 1024; i += 256)
+ {
+ void *partdst = acc_pcopyin (&myblock[i], 256);
+ assert ((uintptr_t) partdst == (uintptr_t) dst + i);
+ }
+ for (i = 0; i < 1024; i += 256)
+ acc_delete (&myblock[i], 256);
+ assert (acc_is_present (myblock, 1024));
+ acc_delete (myblock, 1024);
+ assert (!acc_is_present (myblock, 1024));
+ free (myblock);
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c
new file mode 100644
index 00000000000..d35ab5c4b71
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c
@@ -0,0 +1,35 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <stdlib.h>
+#include <assert.h>
+#include <stdint.h>
+
+int main (int argc, char* argv[])
+{
+ char* block1 = malloc (1024);
+ char *block2 = malloc (1024);
+ char *block3 = malloc (1024);
+ int i;
+ void *dst;
+ for (i = 0; i < 1024; i++)
+ block1[i] = block2[i] = block3[i] = i;
+ #pragma acc data copyin(block1[0:1024]) copyin(block2[0:1024]) \
+ copyin(block3[0:1024])
+ {
+ dst = acc_deviceptr (block2);
+ for (i = 0; i < 1024; i += 256)
+ {
+ void *partdst = acc_pcopyin (&block2[i], 256);
+ assert ((uintptr_t) partdst == (uintptr_t) dst + i);
+ }
+ }
+ assert (acc_is_present (block2, 1024));
+ for (i = 0; i < 1024; i += 256)
+ acc_delete (&block2[i], 256);
+ assert (!acc_is_present (block2, 1024));
+ free (block1);
+ free (block2);
+ free (block3);
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c
index 25ceb3a26af..10d3cbc5cc6 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c
@@ -31,5 +31,5 @@ main (int argc, char **argv)
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] surrounds2 \\\[\[0-9a-fA-FxX\]+,\\\+257\\\]" } */
+/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+257\\\] outside mapped block \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */
/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c
index 65ff440a528..cb32bbcb652 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c
@@ -31,5 +31,3 @@ main (int argc, char **argv)
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] surrounds2 \\\[\[0-9a-fA-FxX\]+,\\\+255\\\]" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c
index fd4dc5971a1..b1f3e71f278 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c
@@ -41,5 +41,5 @@ main (int argc, char **argv)
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] surrounds2 \\\[\[0-9a-fA-FxX\]+,\\\+512\\\]" } */
+/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+512\\\] outside mapped block \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */
/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c
index 9bc9ecc1068..d0e5ffb0691 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c
@@ -28,5 +28,3 @@ main (int argc, char **argv)
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] surrounds2 \\\[\[0-9a-fA-FxX\]+,\\\+254\\\]" } */
-/* { dg-shouldfail "" } */
