OpenACC "present" subarrays: runtime API return value and unmapping fixes

	PR libgomp/92511

	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.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
From-SVN: r278514

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index a7e6870c3eafb63d789430f2769a96af1f882c92..9d4a34c90a3a2840d053670eb5fcc11b5c14e52f 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,20 @@
+2019-11-20  Julian Brown  <julian@codesourcery.com>
+
+       PR libgomp/92511
+
+       * 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.
+
 2019-11-20  Maciej W. Rozycki  <macro@wdc.com>
 
        * testsuite/lib/libgomp.exp (libgomp_init): Add flags to find

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 2f271009fb87689f40a5073a17d31fcbb0297313..aafe88d3a1459580c5f02c87bc45f28f43448e3b 100644 (file)
--- 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))
         {
@@ -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 (file)
index 0000000..7e50f3b
--- /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 = (char *) 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 (file)
index 0000000..00e7da1
--- /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 = (char *) malloc (1024);
+  char *block2 = (char *) malloc (1024);
+  char *block3 = (char *) 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 25ceb3a26af47febe93a8cfe29259d6e3f1ce5a9..10d3cbc5cc6faceb5feb0074d0f3bd0c49410875 100644 (file)
--- 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 65ff440a528cfcc0c0656cfc0b0be0d3e50fabcc..cb32bbcb652e256635acb9dced4d74bfed50fca5 100644 (file)
--- 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 fd4dc5971a197c38bd91f2af79e970fdea142142..b1f3e71f2788d611dfe19fdc460b91e6fb28a3d0 100644 (file)
--- 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 9bc9ecc10689634df77ed63ec255abc83aaaaf26..d0e5ffb069106a42662ed2bd1daecf5034db335f 100644 (file)
--- 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 "" } */