diff mbox series

[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) | expand

Commit Message

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

Comments

Thomas Schwinge July 3, 2020, 3:53 p.m. UTC | #1 
Hi!

On 2020-01-17T13:18:20-0800, Julian Brown <julian@codesourcery.com> wrote:
> 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

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

Such a change -- re-posted as
<http://mid.mail-archive.com/7b2f54faa0a25a7a445ad86bf4726202a1190a4f.1590182783.git.julian@codesourcery.com>,
and
<http://mid.mail-archive.com/0585b4fda1ba5c76dce2ac5053a55e2ceef06041.1592343756.git.julian@codesourcery.com>
-- does fix the 'libgomp.oacc-c-c++-common/pr92843-1.c' test case, but I
don't agree that the change is correct.  Instead, I have pushed
"[OpenACC] Revert always-copyfrom behavior for 'GOMP_MAP_FORCE_FROM' in
'libgomp/oacc-mem.c:goacc_exit_data_internal'",
<http://mid.mail-archive.com/87wo3ky5vn.fsf@euler.schwinge.homeip.net>.
(This moves us past the point where it used to 'abort' before, but
doesn't resolve the XFAIL, yet, which needs changes from "[OpenACC]
Adjust dynamic reference count semantics",
<http://mid.mail-archive.com/5e9472b80dc475214a4a082ef54ee919d7f9dcff.1592343756.git.julian@codesourcery.com>.)


Grüße
 Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
diff mbox series

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