[PATCH 2/3] Don't copy back vars mapped with acc_map_data

[Julian Brown]

This patch prevents "exit data" directives from copying back data that
was mapped with an acc_map_data API call. This matches the behaviour
expected by the pr92843-1.c test, and together with the previous patch
in this series, allows that test to pass (with no other regressions).

Tested alongside other patches in this series with offloading to NVPTX
(with and without the third & final patch).

OK?

Thanks,

Julian

ChangeLog

        PR libgomp/92843

        libgomp/
        * oacc-mem.c (goacc_exit_data_internal): Don't copy-back data mapped
        with acc_map_data on an "exit data" directive.
        * testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Remove XFAIL.  Add
        explanatory comment.
---
 libgomp/oacc-mem.c                                      | 1 +
 libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c | 4 +++-
 2 files changed, 4 insertions(+), 1 deletion(-)

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 45ab2b169d7..783e7f363fb 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1235,6 +1235,7 @@ goacc_exit_data_internal (struct gomp_device_descr 
*acc_dev, size_t mapnum,
              n->refcount--;
 
            if (copyfrom
+               && n->refcount != REFCOUNT_INFINITY
                && (kind != GOMP_MAP_FROM || n->refcount == 0))
              gomp_copy_dev2host (acc_dev, aq, (void *) cur_node.host_start,
                                  (void *) (n->tgt->tgt_start + n->tgt_offset
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c 
b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
index f16c46a37bf..786a12a8504 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
@@ -1,7 +1,6 @@
 /* Verify that 'acc_copyout' etc. is a no-op if there's still a structured
    reference count.  */
 
-/* { dg-xfail-run-if "TODO PR92843" { *-*-* } } */
 /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
 
 #include <assert.h>
@@ -96,6 +95,9 @@ test_acc_map_data ()
   verify_array (h, N, c1);
 
   assign_array (h, N, c1);
+  /* Note that we're not expecting this (nor the copyouts below) to perform
+     an actual "finalize" or copyout since the data was mapped with
+     acc_map_data.  */
 #pragma acc exit data copyout (h[0:N]) finalize
   assert (acc_is_present (h, N));
   verify_array (h, N, c1);
-- 
2.23.0