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

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

[Julian Brown]

[-- Attachment #1: Type: text/plain, Size: 2010 bytes --]

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.

[-- Attachment #2: openacc-present-subarrays-1.diff --]
[-- Type: text/x-patch, Size: 7267 bytes --]

commit 00607b06c8e506b0f0744a230856e1e8776633c3
Author: Julian Brown <julian@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 "" } */

Re: [PATCH] OpenACC "present" subarrays: runtime API return value and unmapping fixes

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 4832 bytes --]

Hi Julian!

On 2019-11-09T01:04:21+0000, Julian Brown <julian@codesourcery.com> wrote:
> 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.

Note that I've filed <https://gcc.gnu.org/PR92511> "[OpenACC] Support
subset subarray mappings", so please reference that one in the
ChangeLog/commit log.

Principal ACK for that problem, and it's solution
('libgomp/oacc-mem.c:present_create_copy' 'if (n)' change).

Then, I was confused, because I couldn't really find wording in the
OpenACC specification that explicitly permits such things.  But given
that, for example, in OpenACC 2.7, 3.2.20. "acc_copyin", 'acc_copyin' is
described to be "equivalent to the 'enter data' directive with a 'copyin'
clause", and the latter supposedly (?) does allow such "subset subarray
mappings", and in 2.7.6. "copyin clause" it is said that "An 'enter data'
directive with a 'copyin' clause is functionally equivalent to a call to
the 'acc_copyin' API routine", that's probably motivation enough to fix
the latter to conform what the former supposedly already is allowing
(though not implementing by means of 'enter data copyin' just calling
'acc_copyin' etc.

I see that 2.7.6. "copyin clause" also states that "The restrictions
regarding subarrays in the present clause apply to this clause", which
per 2.7.4. "present clause" is that "If only a subarray of an array is
present in the current device memory, the 'present' clause must specify
the same subarray, or a subarray that is a proper subset of the subarray
in the data lifetime".  From that we probably are to deduce that it's
fine the other way round (as you've argued): if a subarray of an array
(or, the whole array) is present in the current device memory, the
'present' clause may specify the same subarray, or a subarray that is a
proper subset of the subarray in the data lifetime (my words).  Unless
you object to that, we shall (later) try to get the clarified/amended in
the OpenACC specification.

Indeed I am confirming that such subset subarray mappings do work fine
with PGI 19.4 and 19.10 -- but only when using OpenACC directives, not
necessarily when using OpenACC runtime library calls, huh.  (That's not
our problem to solve, of course, and under the assumption that my test
case has actually been valid.)

Later (not now), we should then also add corresponding testing for actual
'data' etc. constructs being nested in that way.

> 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.

(That's where PGI fails at runtime, but I have not analyzed how exactly
this fails -- let's first clarify that with OpenACC Technical Committee,
later on.)

> 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.

These should then actually be removed, or re-written, because in their
current form they no longer make much sense, as far as I can tell:

For example, 'libgomp.oacc-c-c++-common/lib-22.c':

    acc_copyin (h, N);

... followed by:

    acc_copyout (h + 1, N - 1);

... is now meant to no longer abort with a "surrounds2" message, but
instead we now expect success, and '!acc_is_present'.

I'll take care of that later on -- I have some more tests to add anyway.

> It's still an error to try to copy data beyond
> the bounds of a mapped block, and other existing tests cover those
> cases.

ACK.

> 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.

But please back out this one, for it's not related to this bug fix, and
we shall take care of that in a later patch.  (No need for you to re-post
that one just for this.)

> Tested with offloading to nvptx. OK for trunk?

I'm see C++ compilation failures the new libgomp test cases; OK with
these resolved.  To record the review effort, please include
"Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>" in the commit
log, see <https://gcc.gnu.org/wiki/Reviewed-by>.


Grüße
 Thomas

[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]

Re: [PATCH] OpenACC "present" subarrays: runtime API return value and unmapping fixes

[Thomas Schwinge]

[-- Attachment #1.1: Type: text/plain, Size: 2681 bytes --]

Hi!

On 2019-11-14T17:02:02+0100, I wrote:
> [...] I couldn't really find wording in the
> OpenACC specification that explicitly permits such things.  But given
> that, for example, in OpenACC 2.7, 3.2.20. "acc_copyin", 'acc_copyin' is
> described to be "equivalent to the 'enter data' directive with a 'copyin'
> clause", and the latter supposedly (?) does allow such "subset subarray
> mappings", and in 2.7.6. "copyin clause" it is said that "An 'enter data'
> directive with a 'copyin' clause is functionally equivalent to a call to
> the 'acc_copyin' API routine", that's probably motivation enough to fix
> the latter to conform what the former supposedly already is allowing
> (though not implementing by means of 'enter data copyin' just calling
> 'acc_copyin' etc.
>
> I see that 2.7.6. "copyin clause" also states that "The restrictions
> regarding subarrays in the present clause apply to this clause", which
> per 2.7.4. "present clause" is that "If only a subarray of an array is
> present in the current device memory, the 'present' clause must specify
> the same subarray, or a subarray that is a proper subset of the subarray
> in the data lifetime".  From that we probably are to deduce that it's
> fine the other way round (as you've argued): if a subarray of an array
> (or, the whole array) is present in the current device memory, the
> 'present' clause may specify the same subarray, or a subarray that is a
> proper subset of the subarray in the data lifetime (my words).  Unless
> you object to that, we shall (later) try to get the clarified/amended in
> the OpenACC specification.

I filed <https://github.com/OpenACC/openacc-spec/issues/247> "Subset
subarray restrictions".


> Later (not now), we should then also add corresponding testing for actual
> 'data' etc. constructs being nested in that way.

> On 2019-11-09T01:04:21+0000, Julian Brown <julian@codesourcery.com> wrote:
>> a couple of existing "shouldfail" tests no longer fail, and have been
>> adjusted accordingly.
>
> These should then actually be removed, or re-written, because in their
> current form they no longer make much sense, as far as I can tell:
>
> For example, 'libgomp.oacc-c-c++-common/lib-22.c':
>
>     acc_copyin (h, N);
>
> ... followed by:
>
>     acc_copyout (h + 1, N - 1);
>
> ... is now meant to no longer abort with a "surrounds2" message, but
> instead we now expect success, and '!acc_is_present'.
>
> I'll take care of that later on -- I have some more tests to add anyway.

See attached '[PR92511] More testing for OpenACC "present" subarrays',
committed to trunk in r279122.


Grüße
 Thomas



[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1.2: 0001-PR92511-More-testing-for-OpenACC-present-subar.trunk.patch --]
[-- Type: text/x-diff, Size: 28312 bytes --]

From 2d5187149761bb9566b2c221c9c7ae7a18c92822 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon, 9 Dec 2019 11:40:36 +0000
Subject: [PATCH] [PR92511] More testing for OpenACC "present" subarrays

In particular, "subset subarrays".

	libgomp/
	PR libgomp/92511
	* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c: Remove
	this file...
	* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c: ..., and
	this file...
	* testsuite/libgomp.oacc-c-c++-common/lib-22.c: ..., and this
	file...
	* testsuite/libgomp.oacc-c-c++-common/lib-30.c: ..., and this
	file...
	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-p.c:
	... with their content moved into, and extended in this new file.
	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-a.c:
	New file.
	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-p.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-a.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-2.c:
	Likewise.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@279122 138bc75d-0d04-0410-961f-82ee72b054a4
---
 libgomp/ChangeLog                             |  20 +
 .../copyin-devptr-1.c                         |  28 -
 .../copyin-devptr-2.c                         |  35 --
 .../libgomp.oacc-c-c++-common/lib-22.c        |  33 --
 .../libgomp.oacc-c-c++-common/lib-30.c        |  30 -
 .../subset-subarray-mappings-1-d-a.c          |   7 +
 .../subset-subarray-mappings-1-d-p.c          |   7 +
 .../subset-subarray-mappings-1-r-a.c          |   7 +
 .../subset-subarray-mappings-1-r-p.c          | 514 ++++++++++++++++++
 .../subset-subarray-mappings-2.c              | 115 ++++
 10 files changed, 670 insertions(+), 126 deletions(-)
 delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c
 delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c
 delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c
 delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-a.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-p.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-a.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-p.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-2.c

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 51a00a3a46c..739a76d48ac 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,5 +1,25 @@
 2019-12-09  Thomas Schwinge  <thomas@codesourcery.com>
 
+	PR libgomp/92511
+	* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c: Remove
+	this file...
+	* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c: ..., and
+	this file...
+	* testsuite/libgomp.oacc-c-c++-common/lib-22.c: ..., and this
+	file...
+	* testsuite/libgomp.oacc-c-c++-common/lib-30.c: ..., and this
+	file...
+	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-p.c:
+	... with their content moved into, and extended in this new file.
+	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-a.c:
+	New file.
+	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-p.c:
+	Likewise.
+	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-a.c:
+	Likewise.
+	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-2.c:
+	Likewise.
+
 	* testsuite/libgomp.oacc-c-c++-common/map-data-1.c: New file.
 
 	PR libgomp/92854
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
deleted file mode 100644
index 7e50f3b892e..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c
+++ /dev/null
@@ -1,28 +0,0 @@
-/* { 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
deleted file mode 100644
index 00e7da1f128..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c
+++ /dev/null
@@ -1,35 +0,0 @@
-/* { 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-22.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c
deleted file mode 100644
index cb32bbcb652..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c
+++ /dev/null
@@ -1,33 +0,0 @@
-/* Exercise acc_copyin and acc_copyout on nvidia targets.  */
-
-/* { dg-do run { target openacc_nvidia_accel_selected } } */
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
-  const int N = 256;
-  int i;
-  unsigned char *h;
-
-  h = (unsigned char *) malloc (N);
-
-  for (i = 0; i < N; i++)
-    {
-      h[i] = i;
-    }
-
-  (void) acc_copyin (h, N);
-
-  fprintf (stderr, "CheCKpOInT\n");
-  acc_copyout (h + 1, N - 1);
-
-  free (h);
-
-  return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c
deleted file mode 100644
index d0e5ffb0691..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c
+++ /dev/null
@@ -1,30 +0,0 @@
-/* Exercise an invalid partial acc_delete on nvidia targets.  */
-
-/* { dg-do run { target openacc_nvidia_accel_selected } } */
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
-  const int N = 256;
-  unsigned char *h;
-  void *d;
-
-  h = (unsigned char *) malloc (N);
-
-  d = acc_create (h, N);
-  if (!d)
-    abort ();
-
-  fprintf (stderr, "CheCKpOInT\n");
-  acc_delete (h, N - 2);
-
-  free (h);
-
-  return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-a.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-a.c
new file mode 100644
index 00000000000..1d168c2e585
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-a.c
@@ -0,0 +1,7 @@
+/* Test "subset" subarray mappings
+   { dg-additional-options "-DOPENACC_DIRECTIVES" } using OpenACC directives,
+   { dg-additional-options "-DARRAYS" } using arrays.  */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include "subset-subarray-mappings-1-r-p.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-p.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-p.c
new file mode 100644
index 00000000000..68ed0ce3eca
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-p.c
@@ -0,0 +1,7 @@
+/* Test "subset" subarray mappings
+   { dg-additional-options "-DOPENACC_DIRECTIVES" } using OpenACC directives,
+   { dg-additional-options "-DPOINTERS" } using pointers.  */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include "subset-subarray-mappings-1-r-p.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-a.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-a.c
new file mode 100644
index 00000000000..5c0fd040d87
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-a.c
@@ -0,0 +1,7 @@
+/* Test "subset" subarray mappings
+   { dg-additional-options "-DOPENACC_RUNTIME" } using OpenACC Runtime Library routines,
+   { dg-additional-options "-DARRAYS" } using arrays.  */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include "subset-subarray-mappings-1-r-p.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-p.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-p.c
new file mode 100644
index 00000000000..9b5d83c66dd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-p.c
@@ -0,0 +1,514 @@
+/* Test "subset" subarray mappings
+   { dg-additional-options "-DOPENACC_RUNTIME" } using OpenACC Runtime Library routines,
+   { dg-additional-options "-DPOINTERS" } using pointers.  */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#if OPENACC_RUNTIME
+#elif OPENACC_DIRECTIVES
+#else
+# error
+#endif
+
+#if POINTERS
+#elif ARRAYS
+#else
+# error
+#endif
+
+
+#include <openacc.h>
+#include <acc_prof.h>
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+#include <stdint.h>
+#include <stdbool.h>
+
+
+static bool cb_ev_alloc_expected;
+static size_t cb_ev_alloc_bytes;
+static const void *cb_ev_alloc_device_ptr;
+static void
+cb_ev_alloc (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info)
+{
+  assert (cb_ev_alloc_expected);
+  cb_ev_alloc_expected = false;
+
+  cb_ev_alloc_bytes = event_info->data_event.bytes;
+  cb_ev_alloc_device_ptr = event_info->data_event.device_ptr;
+}
+
+static bool cb_ev_free_expected;
+static const void *cb_ev_free_device_ptr;
+static void
+cb_ev_free (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info)
+{
+  assert (cb_ev_free_expected);
+  cb_ev_free_expected = false;
+
+  cb_ev_free_device_ptr = event_info->data_event.device_ptr;
+}
+
+
+/* Match the alignment processing that
+   'libgomp/target.c:gomp_map_vars_internal' is doing; simplified, not
+   considering special alignment requirements of certain data types.  */
+
+static size_t
+aligned_size (size_t tgt_size)
+{
+  size_t tgt_align = sizeof (void *);
+  return tgt_size + tgt_align - 1;
+}
+
+static const void *
+aligned_address (const void *tgt_start)
+{
+  size_t tgt_align = sizeof (void *);
+  return (void *) (((uintptr_t) tgt_start + tgt_align - 1) & ~(tgt_align - 1));
+}
+
+
+#define SIZE 1024
+#define SUBSET 32
+
+
+static void
+f1 (void)
+{
+  cb_ev_alloc_expected = false;
+  cb_ev_free_expected = false;
+  acc_prof_register (acc_ev_alloc, cb_ev_alloc, acc_reg);
+  acc_prof_register (acc_ev_free, cb_ev_free, acc_reg);
+
+#if POINTERS
+  char* myblock = (char *) malloc (SIZE);
+#else
+  char myblock[SIZE];
+#endif
+  int i;
+  void *dst;
+  for (i = 0; i < SIZE; i++)
+    myblock[i] = i;
+
+  cb_ev_alloc_expected = true;
+#if OPENACC_RUNTIME
+  dst = acc_copyin (myblock, SIZE);
+#else
+# if POINTERS
+#  pragma acc enter data copyin (myblock[0:SIZE])
+# else
+#  pragma acc enter data copyin (myblock)
+# endif
+  dst = acc_deviceptr (myblock);
+#endif
+  assert (dst);
+  assert (!cb_ev_alloc_expected);
+  assert (cb_ev_alloc_bytes == aligned_size (SIZE));
+  assert (aligned_address (cb_ev_alloc_device_ptr) == dst);
+  for (i = 0; i < SIZE; i += SUBSET)
+    {
+      void *partdst = acc_deviceptr (&myblock[i]);
+      assert ((uintptr_t) partdst == (uintptr_t) dst + i);
+      assert (acc_hostptr (partdst) == &myblock[i]);
+    }
+  for (i = 0; i < SIZE; i += SUBSET)
+    {
+      void *partdst;
+#if OPENACC_RUNTIME
+      partdst = acc_pcopyin (&myblock[i], SUBSET);
+#else
+# pragma acc enter data pcopyin (myblock[i:SUBSET])
+      partdst = acc_deviceptr (&myblock[i]);
+#endif
+      assert ((uintptr_t) partdst == (uintptr_t) dst + i);
+    }
+  /* Dereference first half.  */
+  for (i = 0; i < 512; i += SUBSET)
+    {
+      assert (acc_is_present (&myblock[i], SUBSET));
+      assert (acc_is_present (myblock, SIZE));
+#if OPENACC_RUNTIME
+      acc_delete (&myblock[i], SUBSET);
+#else
+# pragma acc exit data delete (myblock[i:SUBSET])
+#endif
+      assert (acc_is_present (&myblock[i], SUBSET));
+      assert (acc_is_present (myblock, SIZE));
+    }
+  /* Dereference all.  */
+#if OPENACC_RUNTIME
+  acc_delete (myblock, SIZE);
+#else
+# if POINTERS
+#  pragma acc exit data delete (myblock[0:SIZE])
+# else
+#  pragma acc exit data delete (myblock)
+# endif
+#endif
+  /* Expect it's still present.  */
+  assert (acc_is_present (myblock, SIZE));
+  /* Dereference second half.  */
+  for (i = 512; i < SIZE; i += SUBSET)
+    {
+      bool last = i >= SIZE - SUBSET;
+
+      assert (acc_is_present (&myblock[i], SUBSET));
+      assert (acc_is_present (myblock, SIZE));
+#if 0 //TODO PR92848
+      if (last)
+	cb_ev_free_expected = true;
+#endif
+#if OPENACC_RUNTIME
+      acc_delete (&myblock[i], SUBSET);
+#else
+# pragma acc exit data delete (myblock[i:SUBSET])
+#endif
+#if 0 //TODO PR92848
+      assert (!cb_ev_free_expected);
+      if (last)
+	assert (cb_ev_free_device_ptr == cb_ev_alloc_device_ptr);
+#endif
+      assert (acc_is_present (&myblock[i], SUBSET) != last);
+      assert (acc_is_present (myblock, SIZE) != last);
+    }
+  /* Expect it's all gone now.  */
+  for (i = 512; i < SIZE; i += SUBSET)
+    assert (!acc_is_present (&myblock[i], SUBSET));
+  assert (!acc_is_present (myblock, SIZE));
+  assert (!acc_is_present (myblock, 1));
+
+#if POINTERS
+  free (myblock);
+#endif
+
+  acc_prof_unregister (acc_ev_alloc, cb_ev_alloc, acc_reg);
+  acc_prof_unregister (acc_ev_free, cb_ev_free, acc_reg);
+}
+
+
+static void
+f2 (void)
+{
+  cb_ev_alloc_expected = false;
+  cb_ev_free_expected = false;
+  acc_prof_register (acc_ev_alloc, cb_ev_alloc, acc_reg);
+  acc_prof_register (acc_ev_free, cb_ev_free, acc_reg);
+
+#if POINTERS
+  char *block1 = (char *) malloc (SIZE);
+  char *block2 = (char *) malloc (SIZE);
+  char *block3 = (char *) malloc (SIZE);
+#else
+  char block1[SIZE];
+  char block2[SIZE];
+  char block3[SIZE];
+#endif
+  int i;
+  for (i = 0; i < SIZE; i++)
+    block1[i] = block2[i] = block3[i] = i;
+
+  cb_ev_alloc_expected = true;
+#if POINTERS
+# pragma acc data copyin(block1[0:SIZE], block2[0:SIZE], block3[0:SIZE])
+#else
+# pragma acc data copyin(block1, block2, block3)
+#endif
+  {
+    void *block1_d = acc_deviceptr (block1);
+    void *block2_d = acc_deviceptr (block2);
+    void *block3_d = acc_deviceptr (block3);
+    assert (!cb_ev_alloc_expected);
+    /* 'block1', 'block2', 'block3' get mapped in one device memory object, in
+       reverse order.  */
+    assert (cb_ev_alloc_bytes == aligned_size (3 * SIZE));
+    assert ((void *) ((uintptr_t) aligned_address (cb_ev_alloc_device_ptr) + 2 * SIZE) == block1_d);
+    assert ((void *) ((uintptr_t) aligned_address (cb_ev_alloc_device_ptr) + 1 * SIZE) == block2_d);
+    assert ((void *) ((uintptr_t) aligned_address (cb_ev_alloc_device_ptr) + 0 * SIZE) == block3_d);
+
+    for (i = 0; i < SIZE; i += SUBSET)
+      {
+	void *block2_part_d;
+#if OPENACC_RUNTIME
+	block2_part_d = acc_pcopyin (&block2[i], SUBSET);
+#else
+# pragma acc enter data pcopyin (block2[i:SUBSET])
+	block2_part_d = acc_deviceptr (&block2[i]);
+#endif
+	assert ((uintptr_t) block2_part_d == (uintptr_t) block2_d + i);
+      }
+  }
+  /* The mappings have been removed, but the device memory object has not yet
+     been 'free'd.  */
+  assert (!acc_is_present (block1, SIZE));
+  assert (acc_is_present (block2, SIZE));
+  assert (!acc_is_present (block3, SIZE));
+  for (i = 0; i < SIZE; i += SUBSET)
+    {
+      bool last = i >= SIZE - SUBSET;
+
+      assert (acc_is_present (block2, SIZE));
+      if (last)
+	cb_ev_free_expected = true;
+#if OPENACC_RUNTIME
+      acc_delete (&block2[i], SUBSET);
+#else
+# pragma acc exit data delete (block2[i:SUBSET])
+#endif
+      assert (!cb_ev_free_expected);
+      if (last)
+	assert (cb_ev_free_device_ptr == cb_ev_alloc_device_ptr);
+    }
+  assert (!acc_is_present (block1, SIZE));
+  assert (!acc_is_present (block2, SIZE));
+  assert (!acc_is_present (block3, SIZE));
+
+#if POINTERS
+  free (block1);
+  free (block2);
+  free (block3);
+#endif
+
+  acc_prof_unregister (acc_ev_alloc, cb_ev_alloc, acc_reg);
+  acc_prof_unregister (acc_ev_free, cb_ev_free, acc_reg);
+}
+
+
+static void
+f3 ()
+{
+  cb_ev_alloc_expected = false;
+  cb_ev_free_expected = false;
+  acc_prof_register (acc_ev_alloc, cb_ev_alloc, acc_reg);
+  acc_prof_register (acc_ev_free, cb_ev_free, acc_reg);
+
+#if POINTERS
+  char *h = (char *) malloc (SIZE);
+#else
+  char h[SIZE];
+#endif
+
+  char *d1;
+  cb_ev_alloc_expected = true;
+#if OPENACC_RUNTIME
+  d1 = (char *) acc_present_or_create (h, SIZE);
+#else
+# if POINTERS
+#  pragma acc enter data present_or_create (h[0:SIZE])
+# else
+#  pragma acc enter data present_or_create (h)
+# endif
+  d1 = (char *) acc_deviceptr (h);
+#endif
+  assert (d1);
+  assert (!cb_ev_alloc_expected);
+  assert (cb_ev_alloc_bytes == aligned_size (SIZE));
+  assert (aligned_address (cb_ev_alloc_device_ptr) == d1);
+  assert (acc_is_present (h, SIZE));
+  assert (acc_is_present (&h[2], SIZE - 2));
+
+  char *d2;
+#if OPENACC_RUNTIME
+  d2 = (char *) acc_present_or_create (&h[2], SIZE - 2);
+#else
+# pragma acc enter data present_or_create (h[2:SIZE - 2])
+  d2 = (char *) acc_deviceptr (&h[2]);
+#endif
+  assert (d2);
+  assert (d1 == d2 - 2);
+  assert (acc_is_present (h, SIZE));
+  assert (acc_is_present (&h[2], SIZE - 2));
+
+  d2 = (char *) acc_deviceptr (&h[2]);
+  assert (d1 == d2 - 2);
+
+#if OPENACC_RUNTIME
+  acc_delete (&h[2], SIZE - 2);
+#else
+# pragma acc exit data delete (h[2:SIZE - 2])
+#endif
+  assert (acc_is_present (h, SIZE));
+  assert (acc_is_present (&h[2], SIZE - 2));
+
+#if 0 //TODO PR92848
+  cb_ev_free_expected = true;
+#endif
+#if OPENACC_RUNTIME
+  acc_delete (h, SIZE);
+#else
+# if POINTERS
+#  pragma acc exit data delete (h[0:SIZE])
+# else
+#  pragma acc exit data delete (h)
+# endif
+#endif
+#if 0 //TODO PR92848
+  assert (!cb_ev_free_expected);
+  assert (cb_ev_free_device_ptr == cb_ev_alloc_device_ptr);
+#endif
+
+  assert (!acc_is_present (h, SIZE));
+  assert (!acc_is_present (&h[2], SIZE - 2));
+  assert (!acc_is_present (h, 1));
+
+# if POINTERS
+  free (h);
+#endif
+
+  acc_prof_unregister (acc_ev_alloc, cb_ev_alloc, acc_reg);
+  acc_prof_unregister (acc_ev_free, cb_ev_free, acc_reg);
+}
+
+
+/* Based on what used to be 'libgomp.oacc-c-c++-common/lib-22.c'.  */
+
+static void
+f_lib_22 (void)
+{
+  cb_ev_alloc_expected = false;
+  cb_ev_free_expected = false;
+  acc_prof_register (acc_ev_alloc, cb_ev_alloc, acc_reg);
+  acc_prof_register (acc_ev_free, cb_ev_free, acc_reg);
+
+  const int c0 = 0;
+  const int c1 = 1;
+
+#if POINTERS
+  char *h = (char *) malloc (SIZE);
+#else
+  char h[SIZE];
+#endif
+
+  memset (h, c0, SIZE);
+  void *d;
+  cb_ev_alloc_expected = true;
+#if OPENACC_RUNTIME
+  d = acc_copyin (h, SIZE);
+#else
+# if POINTERS
+#  pragma acc enter data copyin (h[0:SIZE])
+# else
+#  pragma acc enter data copyin (h)
+# endif
+  d = acc_deviceptr (h);
+#endif
+  assert (d);
+  assert (!cb_ev_alloc_expected);
+  assert (cb_ev_alloc_bytes == aligned_size (SIZE));
+  assert (aligned_address (cb_ev_alloc_device_ptr) == d);
+  /* Overwrite the local memory.  */
+  memset (h, c1, SIZE);
+  /* Now 'copyout' not the whole but only a "subset" subarray, missing one
+     SUBSET at the beginning, and half a SUBSET at the end...  */
+#if 0 //TODO PR92848
+  cb_ev_free_expected = true;
+#endif
+#if OPENACC_RUNTIME
+  acc_copyout (h + SUBSET, SIZE - SUBSET - SUBSET / 2);
+#else
+# pragma acc exit data copyout (h[SUBSET:SIZE - SUBSET - SUBSET / 2])
+#endif
+#if 0 //TODO PR92848
+  /* ..., yet, expect the device memory object to be 'free'd...  */
+  assert (!cb_ev_free_expected);
+  assert (cb_ev_free_device_ptr == cb_ev_alloc_device_ptr);
+#endif
+  /* ..., and the mapping to be removed...  */
+  assert (!acc_is_present (h, SIZE));
+  assert (!acc_is_present (&h[SUBSET], SIZE - SUBSET - SUBSET / 2));
+  assert (!acc_is_present (h, 1));
+  /* ..., but the 'copyout'ed device memory to correspond to just the "subset"
+     subarray.  */
+  for (size_t i = 0; i < SIZE; ++i)
+    {
+      if (i < SUBSET)
+	assert (h[i] == c1);
+      else if (i < SIZE - SUBSET / 2)
+	assert (h[i] == c0);
+      else if (i < SIZE)
+	assert (h[i] == c1);
+    }
+
+#if POINTERS
+  free (h);
+#endif
+
+  acc_prof_unregister (acc_ev_alloc, cb_ev_alloc, acc_reg);
+  acc_prof_unregister (acc_ev_free, cb_ev_free, acc_reg);
+}
+
+
+/* Based on what used to be 'libgomp.oacc-c-c++-common/lib-30.c'.  */
+
+static void
+f_lib_30 (void)
+{
+  cb_ev_alloc_expected = false;
+  cb_ev_free_expected = false;
+  acc_prof_register (acc_ev_alloc, cb_ev_alloc, acc_reg);
+  acc_prof_register (acc_ev_free, cb_ev_free, acc_reg);
+
+#if POINTERS
+  char *h = (char *) malloc (SIZE);
+#else
+  char h[SIZE];
+#endif
+  memset (h, 0, SIZE);
+
+  void *d;
+  cb_ev_alloc_expected = true;
+#if OPENACC_RUNTIME
+  d = acc_create (h, SIZE);
+#else
+# if POINTERS
+#  pragma acc enter data create (h[0:SIZE])
+# else
+#  pragma acc enter data create (h)
+# endif
+  d = acc_deviceptr (h);
+#endif
+  assert (d);
+  assert (!cb_ev_alloc_expected);
+  assert (cb_ev_alloc_bytes == aligned_size (SIZE));
+  assert (aligned_address (cb_ev_alloc_device_ptr) == d);
+
+  /* We 'delete' not the whole but only a "subset" subarray...  */
+#if 0 //TODO PR92848
+  cb_ev_free_expected = true;
+#endif
+#if OPENACC_RUNTIME
+  acc_delete (h, SIZE - SUBSET);
+#else
+# pragma acc exit data delete (h[0:SIZE - SUBSET])
+#endif
+#if 0 //TODO PR92848
+  /* ..., yet, expect the device memory object to be 'free'd...  */
+  assert (!cb_ev_free_expected);
+  assert (cb_ev_free_device_ptr == cb_ev_alloc_device_ptr);
+#endif
+  /* ..., and the mapping to be removed.  */
+  assert (!acc_is_present (h, SIZE));
+  assert (!acc_is_present (h, SIZE - SUBSET));
+  assert (!acc_is_present (h, 1));
+
+#if POINTERS
+  free (h);
+#endif
+
+  acc_prof_unregister (acc_ev_alloc, cb_ev_alloc, acc_reg);
+  acc_prof_unregister (acc_ev_free, cb_ev_free, acc_reg);
+}
+
+
+int
+main ()
+{
+  f1 ();
+  f2 ();
+  f3 ();
+  f_lib_22 ();
+  f_lib_30 ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-2.c
new file mode 100644
index 00000000000..f4e18fa97a7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-2.c
@@ -0,0 +1,115 @@
+/* Test "subset" subarray mappings.  */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <acc_prof.h>
+#include <stdbool.h>
+#include <stdint.h>
+#include <stdlib.h>
+#include <assert.h>
+
+
+static bool cb_ev_alloc_expected;
+static size_t cb_ev_alloc_bytes;
+static const void *cb_ev_alloc_device_ptr;
+static void
+cb_ev_alloc (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info)
+{
+  assert (cb_ev_alloc_expected);
+  cb_ev_alloc_expected = false;
+
+  cb_ev_alloc_bytes = event_info->data_event.bytes;
+  cb_ev_alloc_device_ptr = event_info->data_event.device_ptr;
+}
+
+static bool cb_ev_free_expected;
+static const void *cb_ev_free_device_ptr;
+static void
+cb_ev_free (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info)
+{
+  assert (cb_ev_free_expected);
+  cb_ev_free_expected = false;
+
+  cb_ev_free_device_ptr = event_info->data_event.device_ptr;
+}
+
+
+/* Match the alignment processing that
+   'libgomp/target.c:gomp_map_vars_internal' is doing; simplified, not
+   considering special alignment requirements of certain data types.  */
+
+static size_t
+aligned_size (size_t tgt_size)
+{
+  size_t tgt_align = sizeof (void *);
+  return tgt_size + tgt_align - 1;
+}
+
+static const void *
+aligned_address (const void *tgt_start)
+{
+  size_t tgt_align = sizeof (void *);
+  return (void *) (((uintptr_t) tgt_start + tgt_align - 1) & ~(tgt_align - 1));
+}
+
+
+#define SIZE 1024
+
+
+int
+main ()
+{
+  cb_ev_alloc_expected = false;
+  cb_ev_free_expected = false;
+  acc_prof_register (acc_ev_alloc, cb_ev_alloc, acc_reg);
+  acc_prof_register (acc_ev_free, cb_ev_free, acc_reg);
+
+  char *block1 = (char *) malloc (SIZE);
+  char *block2 = (char *) malloc (SIZE);
+  char *block3 = (char *) malloc (SIZE);
+  cb_ev_alloc_expected = true;
+#pragma acc data create (block1[0:SIZE], block2[0:SIZE], block3[0:SIZE])
+  {
+    void *s_block1_d = acc_deviceptr (&block1[1]);
+    void *s_block2_d = acc_deviceptr (&block2[20]);
+    void *s_block3_d = acc_deviceptr (&block3[300]);
+    assert (!cb_ev_alloc_expected);
+    /* 'block1', 'block2', 'block3' get mapped in one device memory object, in
+       reverse order.  */
+    assert (cb_ev_alloc_bytes == aligned_size (3 * SIZE));
+    assert ((void *) ((uintptr_t) aligned_address (cb_ev_alloc_device_ptr) + 2 * SIZE + 1) == s_block1_d);
+    assert ((void *) ((uintptr_t) aligned_address (cb_ev_alloc_device_ptr) + 1 * SIZE + 20) == s_block2_d);
+    assert ((void *) ((uintptr_t) aligned_address (cb_ev_alloc_device_ptr) + 0 * SIZE + 300) == s_block3_d);
+
+    void *s_block1_p_d = acc_pcopyin (&block1[1], SIZE - 3);
+    void *s_block2_p_d = acc_pcopyin (&block2[20], SIZE - 33);
+    void *s_block3_p_d = acc_pcopyin (&block3[300], SIZE - 333);
+    assert (s_block1_p_d == s_block1_d);
+    assert (s_block2_p_d == s_block2_d);
+    assert (s_block3_p_d == s_block3_d);
+
+    acc_delete (block1, SIZE);
+    acc_delete (block2, SIZE);
+    acc_delete (block3, SIZE);
+    assert (acc_is_present (block1, SIZE));
+    assert (acc_is_present (block2, SIZE));
+    assert (acc_is_present (block3, SIZE));
+
+    cb_ev_free_expected = true;
+  }
+  assert (!cb_ev_free_expected);
+  assert (cb_ev_free_device_ptr == cb_ev_alloc_device_ptr);
+  assert (!acc_is_present (block1, SIZE));
+  assert (!acc_is_present (block2, SIZE));
+  assert (!acc_is_present (block3, SIZE));
+
+  free (block1);
+  free (block2);
+  free (block3);
+
+  acc_prof_unregister (acc_ev_alloc, cb_ev_alloc, acc_reg);
+  acc_prof_unregister (acc_ev_free, cb_ev_free, acc_reg);
+
+  return 0;
+}
-- 
2.17.1


[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]