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

6 files changed, 65 insertions, 6 deletions

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 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
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 25ceb3a..10d3cbc 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 65ff440..cb32bbc 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 fd4dc59..b1f3e71 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 9bc9ecc..d0e5ffb 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 "" } */