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

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

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)

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;
+}

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;
+}

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 "" } */

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 "" } */

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 "" } */

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 "" } */