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

Message ID 4bbb4970cd48424873573a50e627786cc9cf3378.1579292772.git.julian@codesourcery.com
State New
Headers show
Series
  • Mixed static/dynamic data lifetimes with OpenACC (PR92843)
Related show

Commit Message

Julian Brown Jan. 17, 2020, 9:18 p.m.
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(-)

-- 
2.23.0

Patch

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);